From d848704b35318c2498b4950ddd82e341c1dd28ca Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 23 Jun 2014 13:55:09 -0700 Subject: [PATCH 01/19] cuda::DisparityBilateralFilter no longer uses constant memory for parameters Now multiple filters can be used in the same context without stepping on each other. --- .../src/cuda/disparity_bilateral_filter.cu | 44 +++++-------------- .../src/disparity_bilateral_filter.cpp | 8 ++-- 2 files changed, 15 insertions(+), 37 deletions(-) diff --git a/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu b/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu index b5de989ae7..5b16f8c108 100644 --- a/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu +++ b/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu @@ -49,30 +49,6 @@ namespace cv { namespace cuda { namespace device { namespace disp_bilateral_filter { - __constant__ float* ctable_color; - __constant__ float* ctable_space; - __constant__ size_t ctable_space_step; - - __constant__ int cndisp; - __constant__ int cradius; - - __constant__ short cedge_disc; - __constant__ short cmax_disc; - - void disp_load_constants(float* table_color, PtrStepSzf table_space, int ndisp, int radius, short edge_disc, short max_disc) - { - cudaSafeCall( cudaMemcpyToSymbol(ctable_color, &table_color, sizeof(table_color)) ); - cudaSafeCall( cudaMemcpyToSymbol(ctable_space, &table_space.data, sizeof(table_space.data)) ); - size_t table_space_step = table_space.step / sizeof(float); - cudaSafeCall( cudaMemcpyToSymbol(ctable_space_step, &table_space_step, sizeof(size_t)) ); - - cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(cradius, &radius, sizeof(int)) ); - - cudaSafeCall( cudaMemcpyToSymbol(cedge_disc, &edge_disc, sizeof(short)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmax_disc, &max_disc, sizeof(short)) ); - } - template struct DistRgbMax { @@ -95,7 +71,11 @@ namespace cv { namespace cuda { namespace device }; template - __global__ void disp_bilateral_filter(int t, T* disp, size_t disp_step, const uchar* img, size_t img_step, int h, int w) + __global__ void disp_bilateral_filter(int t, T* disp, size_t disp_step, + const uchar* img, size_t img_step, int h, int w, + const float* ctable_color, const float * ctable_space, size_t ctable_space_step, + int cradius, + short cedge_disc, short cmax_disc) { const int y = blockIdx.y * blockDim.y + threadIdx.y; const int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + t) & 1); @@ -178,7 +158,7 @@ namespace cv { namespace cuda { namespace device } template - void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream) + void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *table_color, const float* table_space, size_t table_step, int radius, short edge_disc, short max_disc, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -190,20 +170,20 @@ namespace cv { namespace cuda { namespace device case 1: for (int i = 0; i < iters; ++i) { - disp_bilateral_filter<1><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); + disp_bilateral_filter<1><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols, table_color, table_space, table_step, radius, edge_disc, max_disc); cudaSafeCall( cudaGetLastError() ); - disp_bilateral_filter<1><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); + disp_bilateral_filter<1><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols, table_color, table_space, table_step, radius, edge_disc, max_disc); cudaSafeCall( cudaGetLastError() ); } break; case 3: for (int i = 0; i < iters; ++i) { - disp_bilateral_filter<3><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); + disp_bilateral_filter<3><<>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols, table_color, table_space, table_step, radius, edge_disc, max_disc); cudaSafeCall( cudaGetLastError() ); - disp_bilateral_filter<3><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols); + disp_bilateral_filter<3><<>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols, table_color, table_space, table_step, radius, edge_disc, max_disc); cudaSafeCall( cudaGetLastError() ); } break; @@ -215,8 +195,8 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); - template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); + template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *table_color, const float *table_space, size_t table_step, int radius, short, short, cudaStream_t stream); + template void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *table_color, const float *table_space, size_t table_step, int radius, short, short, cudaStream_t stream); } // namespace bilateral_filter }}} // namespace cv { namespace cuda { namespace cudev diff --git a/modules/cudastereo/src/disparity_bilateral_filter.cpp b/modules/cudastereo/src/disparity_bilateral_filter.cpp index 75cbce48a9..746d5d49cd 100644 --- a/modules/cudastereo/src/disparity_bilateral_filter.cpp +++ b/modules/cudastereo/src/disparity_bilateral_filter.cpp @@ -55,10 +55,8 @@ namespace cv { namespace cuda { namespace device { namespace disp_bilateral_filter { - void disp_load_constants(float* table_color, PtrStepSzf table_space, int ndisp, int radius, short edge_disc, short max_disc); - template - void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream); + void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream); } }}} @@ -165,7 +163,7 @@ namespace const short edge_disc = std::max(short(1), short(ndisp * edge_threshold + 0.5)); const short max_disc = short(ndisp * max_disc_threshold + 0.5); - disp_load_constants(table_color.ptr(), table_space, ndisp, radius, edge_disc, max_disc); + size_t table_space_step = table_space.step / sizeof(float); _dst.create(disp.size(), disp.type()); GpuMat dst = _dst.getGpuMat(); @@ -173,7 +171,7 @@ namespace if (dst.data != disp.data) disp.copyTo(dst, stream); - disp_bilateral_filter(dst, img, img.channels(), iters, StreamAccessor::getStream(stream)); + disp_bilateral_filter(dst, img, img.channels(), iters, table_color.ptr(), (float *)table_space.data, table_space_step, radius, edge_disc, max_disc, StreamAccessor::getStream(stream)); } void DispBilateralFilterImpl::apply(InputArray _disp, InputArray _image, OutputArray dst, Stream& stream) From 1533d0448f130cbec92c9bba0e200f9864824d2f Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Tue, 24 Jun 2014 11:09:46 -0700 Subject: [PATCH 02/19] cdata_weight always positive; fewer multiplications. --- modules/cudastereo/src/cuda/stereocsbp.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index b1426607dd..b10007e003 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -103,7 +103,7 @@ namespace cv { namespace cuda { namespace device { static __device__ __forceinline__ float compute(const uchar* left, const uchar* right) { - return fmin(cdata_weight * ::abs((int)*left - *right), cdata_weight * cmax_data_term); + return cdata_weight * fmin( ::abs((int)*left - *right), cmax_data_term); } }; template <> struct DataCostPerPixel<3> @@ -114,7 +114,7 @@ namespace cv { namespace cuda { namespace device float tg = 0.587f * ::abs((int)left[1] - right[1]); float tr = 0.299f * ::abs((int)left[2] - right[2]); - return fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term); + return cdata_weight * fmin(tr + tg + tb, cmax_data_term); } }; template <> struct DataCostPerPixel<4> @@ -128,7 +128,7 @@ namespace cv { namespace cuda { namespace device float tg = 0.587f * ::abs((int)l.y - r.y); float tr = 0.299f * ::abs((int)l.z - r.z); - return fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term); + return cdata_weight * fmin(tr + tg + tb, cmax_data_term); } }; From e532bd50d5f631cb27585176087096461012073d Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Thu, 26 Jun 2014 12:06:25 -0700 Subject: [PATCH 03/19] Move shared interface to header file rather than repeating. --- .../cudastereo/src/cuda/disparity_bilateral_filter.cu | 2 ++ .../cudastereo/src/cuda/disparity_bilateral_filter.hpp | 8 ++++++++ modules/cudastereo/src/disparity_bilateral_filter.cpp | 9 +-------- 3 files changed, 11 insertions(+), 8 deletions(-) create mode 100644 modules/cudastereo/src/cuda/disparity_bilateral_filter.hpp diff --git a/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu b/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu index 5b16f8c108..a9f2d2650c 100644 --- a/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu +++ b/modules/cudastereo/src/cuda/disparity_bilateral_filter.cu @@ -45,6 +45,8 @@ #include "opencv2/core/cuda/common.hpp" #include "opencv2/core/cuda/limits.hpp" +#include "cuda/disparity_bilateral_filter.hpp" + namespace cv { namespace cuda { namespace device { namespace disp_bilateral_filter diff --git a/modules/cudastereo/src/cuda/disparity_bilateral_filter.hpp b/modules/cudastereo/src/cuda/disparity_bilateral_filter.hpp new file mode 100644 index 0000000000..95be834573 --- /dev/null +++ b/modules/cudastereo/src/cuda/disparity_bilateral_filter.hpp @@ -0,0 +1,8 @@ +namespace cv { namespace cuda { namespace device +{ + namespace disp_bilateral_filter + { + template + void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream); + } +}}} diff --git a/modules/cudastereo/src/disparity_bilateral_filter.cpp b/modules/cudastereo/src/disparity_bilateral_filter.cpp index 746d5d49cd..c59e3b2cb4 100644 --- a/modules/cudastereo/src/disparity_bilateral_filter.cpp +++ b/modules/cudastereo/src/disparity_bilateral_filter.cpp @@ -51,14 +51,7 @@ Ptr cv::cuda::createDisparityBilateralFilter(int #else /* !defined (HAVE_CUDA) */ -namespace cv { namespace cuda { namespace device -{ - namespace disp_bilateral_filter - { - template - void disp_bilateral_filter(PtrStepSz disp, PtrStepSzb img, int channels, int iters, const float *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream); - } -}}} +#include "cuda/disparity_bilateral_filter.hpp" namespace { From fe29ed461cb61f9e0105ac18a2ef8aa727156ca4 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Thu, 26 Jun 2014 14:06:45 -0700 Subject: [PATCH 04/19] Move shared interface to header file rather than repeating. --- modules/cudastereo/src/cuda/stereocsbp.cu | 2 ++ modules/cudastereo/src/cuda/stereocsbp.hpp | 31 +++++++++++++++++++++ modules/cudastereo/src/stereocsbp.cpp | 32 +--------------------- 3 files changed, 34 insertions(+), 31 deletions(-) create mode 100644 modules/cudastereo/src/cuda/stereocsbp.hpp diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index b10007e003..dc3dbcce9e 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -48,6 +48,8 @@ #include "opencv2/core/cuda/reduce.hpp" #include "opencv2/core/cuda/functional.hpp" +#include "cuda/stereocsbp.hpp" + namespace cv { namespace cuda { namespace device { namespace stereocsbp diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp new file mode 100644 index 0000000000..9b90732ffb --- /dev/null +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -0,0 +1,31 @@ +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, + const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& temp); + + template + void init_data_cost(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 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 + void init_message(T* u_new, T* d_new, T* l_new, T* r_new, + const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, + T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, + T* data_cost_selected, const T* data_cost, size_t msg_step, + int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream); + + template + void calc_all_iterations(T* u, T* d, T* l, T* r, const T* data_cost_selected, + const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); + + template + void compute_disp(const T* u, const T* d, const T* l, const T* r, const T* data_cost_selected, const T* disp_selected, size_t msg_step, + const PtrStepSz& disp, int nr_plane, cudaStream_t stream); + } +}}} diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 474562baf2..2d25bee37d 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -53,37 +53,7 @@ Ptr cv::cuda::createStereoConstantSpaceBP(int, int, #else /* !defined (HAVE_CUDA) */ -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, - const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& temp); - - template - void init_data_cost(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 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 - void init_message(T* u_new, T* d_new, T* l_new, T* r_new, - const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, - T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, - T* data_cost_selected, const T* data_cost, size_t msg_step, - int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream); - - template - void calc_all_iterations(T* u, T* d, T* l, T* r, const T* data_cost_selected, - const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); - - template - void compute_disp(const T* u, const T* d, const T* l, const T* r, const T* data_cost_selected, const T* disp_selected, size_t msg_step, - const PtrStepSz& disp, int nr_plane, cudaStream_t stream); - } -}}} +#include "cuda/stereocsbp.hpp" namespace { From 2982e77495356985e4dc8c4e453f2e2b399a1388 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Thu, 26 Jun 2014 15:31:21 -0700 Subject: [PATCH 05/19] Pass in images and scratch space so that multiple copies can run concurrently. --- modules/cudastereo/src/cuda/stereocsbp.cu | 100 ++++++++++----------- modules/cudastereo/src/cuda/stereocsbp.hpp | 10 +-- modules/cudastereo/src/stereocsbp.cpp | 20 +++-- 3 files changed, 62 insertions(+), 68 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index dc3dbcce9e..a8c70c92de 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -72,13 +72,9 @@ namespace cv { namespace cuda { namespace device __constant__ size_t cdisp_step1; __constant__ size_t cdisp_step2; - __constant__ uchar* cleft; - __constant__ uchar* cright; - __constant__ uchar* ctemp; - void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th, - const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& temp) + size_t leftstep) { cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); @@ -89,11 +85,7 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(cimg_step, &left.step, sizeof(size_t)) ); - - cudaSafeCall( cudaMemcpyToSymbol(cleft, &left.data, sizeof(left.data)) ); - cudaSafeCall( cudaMemcpyToSymbol(cright, &right.data, sizeof(right.data)) ); - cudaSafeCall( cudaMemcpyToSymbol(ctemp, &temp.data, sizeof(temp.data)) ); + cudaSafeCall( cudaMemcpyToSymbol(cimg_step, &leftstep, sizeof(size_t)) ); } /////////////////////////////////////////////////////////////// @@ -135,7 +127,7 @@ namespace cv { namespace cuda { namespace device }; template - __global__ void get_first_k_initial_global(T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane) + __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -169,7 +161,7 @@ namespace cv { namespace cuda { namespace device template - __global__ void get_first_k_initial_local(T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane) + __global__ void get_first_k_initial_local(uchar *ctemp, T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -225,7 +217,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(int h, int w, int level) + __global__ void init_data_cost(uchar *cleft, uchar *cright, uchar *ctemp, int h, int w, int level) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -265,7 +257,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost_reduce(int level, int rows, int cols, int h) + __global__ void init_data_cost_reduce(uchar *cleft, uchar *cright, uchar *ctemp, int level, int rows, int cols, int h) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -313,7 +305,7 @@ namespace cv { namespace cuda { namespace device template - void init_data_cost_caller_(int /*rows*/, int /*cols*/, int h, int w, int level, int /*ndisp*/, int channels, cudaStream_t stream) + void init_data_cost_caller_(uchar *cleft, uchar *cright, uchar *ctemp, 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); @@ -323,15 +315,15 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost<<>>(h, w, level); break; - case 3: init_data_cost<<>>(h, w, level); break; - case 4: init_data_cost<<>>(h, w, level); break; + 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; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void init_data_cost_reduce_caller_(int rows, int cols, int h, int w, int level, int ndisp, int channels, cudaStream_t stream) + void init_data_cost_reduce_caller_(uchar *cleft, uchar *cright, uchar *ctemp, 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); @@ -342,19 +334,19 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost_reduce<<>>(level, rows, cols, h); break; - case 3: init_data_cost_reduce<<>>(level, rows, cols, h); break; - case 4: init_data_cost_reduce<<>>(level, rows, cols, h); break; + 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; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void init_data_cost(int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, + void init_data_cost(uchar *cleft, uchar *cright, uchar *ctemp, 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)(int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream); + typedef void (*InitDataCostCaller)(uchar *cleft, uchar *cright, uchar *ctemp, int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -367,7 +359,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](rows, cols, h, w, level, ndisp, channels, stream); + init_data_cost_callers[level](cleft, cright, ctemp, rows, cols, h, w, level, ndisp, channels, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -380,9 +372,9 @@ namespace cv { namespace cuda { namespace device grid.y = divUp(h, threads.y); if (use_local_init_data_cost == true) - get_first_k_initial_local<<>> (data_cost_selected, disp_selected_pyr, h, w, nr_plane); + get_first_k_initial_local<<>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane); else - get_first_k_initial_global<<>>(data_cost_selected, disp_selected_pyr, h, w, nr_plane); + get_first_k_initial_global<<>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane); cudaSafeCall( cudaGetLastError() ); @@ -390,10 +382,10 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - template void init_data_cost(int rows, int cols, short* disp_selected_pyr, short* data_cost_selected, size_t msg_step, + template void init_data_cost(uchar *cleft, uchar *cright, uchar *ctemp, 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(int rows, int cols, float* disp_selected_pyr, float* data_cost_selected, size_t msg_step, + template void init_data_cost(uchar *cleft, uchar *cright, uchar *ctemp, 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); /////////////////////////////////////////////////////////////// @@ -401,7 +393,7 @@ namespace cv { namespace cuda { namespace device /////////////////////////////////////////////////////////////// template - __global__ void compute_data_cost(const T* selected_disp_pyr, T* data_cost_, int h, int w, int level, int nr_plane) + __global__ void compute_data_cost(uchar *cleft, uchar *cright, 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; @@ -444,7 +436,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_data_cost_reduce(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(uchar *cleft, uchar *cright, 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; @@ -494,7 +486,7 @@ namespace cv { namespace cuda { namespace device } template - void compute_data_cost_caller_(const T* disp_selected_pyr, T* data_cost, int /*rows*/, int /*cols*/, + void compute_data_cost_caller_(uchar *cleft, uchar *cright, 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); @@ -505,15 +497,15 @@ namespace cv { namespace cuda { namespace device switch(channels) { - case 1: compute_data_cost<<>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break; - case 3: compute_data_cost<<>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break; - case 4: compute_data_cost<<>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break; + 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; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void compute_data_cost_reduce_caller_(const T* disp_selected_pyr, T* data_cost, int rows, int cols, + void compute_data_cost_reduce_caller_(uchar *cleft, uchar *cright, 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; @@ -525,18 +517,18 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: compute_data_cost_reduce<<>>(disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; - case 3: compute_data_cost_reduce<<>>(disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; - case 4: compute_data_cost_reduce<<>>(disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; + 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; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template - void compute_data_cost(const T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(uchar *cleft, uchar *cright, 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 T* disp_selected_pyr, T* data_cost, int rows, int cols, + typedef void (*ComputeDataCostCaller)(uchar *cleft, uchar *cright, 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[] = @@ -552,17 +544,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](disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, stream); + callers[level](cleft, cright, 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 short* disp_selected_pyr, short* data_cost, size_t msg_step, + template void compute_data_cost(uchar *cleft, uchar *cright, 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 float* disp_selected_pyr, float* data_cost, size_t msg_step, + template void compute_data_cost(uchar *cleft, uchar *cright, 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); @@ -605,7 +597,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_message(T* u_new_, T* d_new_, T* l_new_, T* r_new_, + __global__ void init_message(uchar *ctemp, T* u_new_, T* d_new_, T* l_new_, T* r_new_, const T* u_cur_, const T* d_cur_, const T* l_cur_, const T* r_cur_, T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, T* data_cost_selected_, const T* data_cost_, @@ -655,7 +647,7 @@ namespace cv { namespace cuda { namespace device template - void init_message(T* u_new, T* d_new, T* l_new, T* r_new, + void init_message(uchar *ctemp, T* u_new, T* d_new, T* l_new, T* r_new, const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, T* data_cost_selected, const T* data_cost, size_t msg_step, @@ -674,7 +666,7 @@ namespace cv { namespace cuda { namespace device grid.x = divUp(w, threads.x); grid.y = divUp(h, threads.y); - init_message<<>>(u_new, d_new, l_new, r_new, + init_message<<>>(ctemp, u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur, selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, @@ -686,13 +678,13 @@ namespace cv { namespace cuda { namespace device } - template void init_message(short* u_new, short* d_new, short* l_new, short* r_new, + template void init_message(uchar *ctemp, short* u_new, short* d_new, short* l_new, short* r_new, const short* u_cur, const short* d_cur, const short* l_cur, const short* r_cur, short* selected_disp_pyr_new, const short* selected_disp_pyr_cur, short* data_cost_selected, const short* data_cost, size_t msg_step, int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream); - template void init_message(float* u_new, float* d_new, float* l_new, float* r_new, + template void init_message(uchar *ctemp, float* u_new, float* d_new, float* l_new, float* r_new, const float* u_cur, const float* d_cur, const float* l_cur, const float* r_cur, float* selected_disp_pyr_new, const float* selected_disp_pyr_cur, float* data_cost_selected, const float* data_cost, size_t msg_step, @@ -738,7 +730,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_message(T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i) + __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i) { int y = blockIdx.y * blockDim.y + threadIdx.y; int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + i) & 1); @@ -765,7 +757,7 @@ namespace cv { namespace cuda { namespace device template - void calc_all_iterations(T* u, T* d, T* l, T* r, const T* data_cost_selected, + void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream) { size_t disp_step = msg_step * h; @@ -780,17 +772,17 @@ namespace cv { namespace cuda { namespace device for(int t = 0; t < iters; ++t) { - compute_message<<>>(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1); + compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1); cudaSafeCall( cudaGetLastError() ); } if (stream == 0) cudaSafeCall( cudaDeviceSynchronize() ); }; - template void calc_all_iterations(short* u, short* d, short* l, short* r, const short* data_cost_selected, const short* selected_disp_pyr_cur, size_t msg_step, + template void calc_all_iterations(uchar *ctemp, short* u, short* d, short* l, short* r, const short* data_cost_selected, const short* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); - template void calc_all_iterations(float* u, float* d, float* l, float* r, const float* data_cost_selected, const float* selected_disp_pyr_cur, size_t msg_step, + template void calc_all_iterations(uchar *ctemp, float* u, float* d, float* l, float* r, const float* data_cost_selected, const float* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 9b90732ffb..ff5e0e1f0c 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -3,25 +3,25 @@ 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, - const PtrStepSzb& left, const PtrStepSzb& right, const PtrStepSzb& temp); + size_t leftstep); template - void init_data_cost(int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, + void init_data_cost(uchar *left, uchar *right, uchar *ctemp, 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 T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(uchar *left, uchar *right, 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 - void init_message(T* u_new, T* d_new, T* l_new, T* r_new, + void init_message(uchar *ctemp, T* u_new, T* d_new, T* l_new, T* r_new, const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, T* data_cost_selected, const T* data_cost, size_t msg_step, int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream); template - void calc_all_iterations(T* u, T* d, T* l, T* r, const T* data_cost_selected, + void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); template diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 2d25bee37d..55afd8c6b9 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, right, temp_); + load_constants(ndisp_, max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_, left.step); l[0].setTo(0, _stream); d[0].setTo(0, _stream); @@ -245,17 +245,18 @@ namespace { if (i == levels_ - 1) { - init_data_cost(left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), + init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), 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(disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, + compute_data_cost(left.ptr(), right.ptr(), 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; - init_message(u[new_idx].ptr(), d[new_idx].ptr(), l[new_idx].ptr(), r[new_idx].ptr(), + init_message(temp_.ptr(), + u[new_idx].ptr(), d[new_idx].ptr(), l[new_idx].ptr(), r[new_idx].ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), disp_selected_pyr[new_idx].ptr(), disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), data_cost.ptr(), elem_step, rows_pyr[i], @@ -264,7 +265,7 @@ namespace cur_idx = new_idx; } - calc_all_iterations(u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), + calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, stream); } @@ -275,17 +276,18 @@ namespace { if (i == levels_ - 1) { - init_data_cost(left.rows, left.cols, disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), + init_data_cost(left.ptr(), right.ptr(), temp_.ptr(), 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(disp_selected_pyr[cur_idx].ptr(), data_cost.ptr(), elem_step, + compute_data_cost(left.ptr(), right.ptr(), 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; - init_message(u[new_idx].ptr(), d[new_idx].ptr(), l[new_idx].ptr(), r[new_idx].ptr(), + init_message(temp_.ptr(), + u[new_idx].ptr(), d[new_idx].ptr(), l[new_idx].ptr(), r[new_idx].ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), disp_selected_pyr[new_idx].ptr(), disp_selected_pyr[cur_idx].ptr(), data_cost_selected.ptr(), data_cost.ptr(), elem_step, rows_pyr[i], @@ -294,7 +296,7 @@ namespace cur_idx = new_idx; } - calc_all_iterations(u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), + calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, stream); } From d8d946a45838ab744f2f406d932f9f5ab1635d06 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Thu, 26 Jun 2014 15:58:05 -0700 Subject: [PATCH 06/19] Constify cuda csbp --- modules/cudastereo/src/cuda/stereocsbp.cu | 32 +++++++++++----------- modules/cudastereo/src/cuda/stereocsbp.hpp | 4 +-- 2 files changed, 18 insertions(+), 18 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index a8c70c92de..fa17742349 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -217,7 +217,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(uchar *cleft, uchar *cright, uchar *ctemp, int h, int w, int level) + __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, int h, int w, int level) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -257,7 +257,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost_reduce(uchar *cleft, 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, int level, int rows, int cols, int h) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -305,7 +305,7 @@ namespace cv { namespace cuda { namespace device template - void init_data_cost_caller_(uchar *cleft, 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, 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); @@ -323,7 +323,7 @@ namespace cv { namespace cuda { namespace device } template - void init_data_cost_reduce_caller_(uchar *cleft, 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, 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); @@ -342,11 +342,11 @@ namespace cv { namespace cuda { namespace device } template - void init_data_cost(uchar *cleft, 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, 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)(uchar *cleft, 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, int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -382,10 +382,10 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - template void init_data_cost(uchar *cleft, 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, 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(uchar *cleft, 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, 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 +393,7 @@ namespace cv { namespace cuda { namespace device /////////////////////////////////////////////////////////////// template - __global__ void compute_data_cost(uchar *cleft, 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, 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 +436,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_data_cost_reduce(uchar *cleft, 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, 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 +486,7 @@ namespace cv { namespace cuda { namespace device } template - void compute_data_cost_caller_(uchar *cleft, 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, 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); @@ -505,7 +505,7 @@ namespace cv { namespace cuda { namespace device } template - void compute_data_cost_reduce_caller_(uchar *cleft, 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, 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; @@ -525,10 +525,10 @@ namespace cv { namespace cuda { namespace device } template - void compute_data_cost(uchar *cleft, uchar *cright, const T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(const uchar *cleft, const uchar *cright, 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)(uchar *cleft, uchar *cright, const T* disp_selected_pyr, T* data_cost, int rows, int cols, + typedef void (*ComputeDataCostCaller)(const uchar *cleft, const uchar *cright, 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[] = @@ -551,10 +551,10 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - template void compute_data_cost(uchar *cleft, 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, 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(uchar *cleft, 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, 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 ff5e0e1f0c..72b5ab04c7 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -6,11 +6,11 @@ namespace cv { namespace cuda { namespace device size_t leftstep); template - void init_data_cost(uchar *left, 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, 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(uchar *left, uchar *right, const T* disp_selected_pyr, T* data_cost, size_t msg_step, + void compute_data_cost(const uchar *left, const uchar *right, 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 From 2832cfdfe581c8a18c695d6737d8ab92f7f6f846 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Fri, 27 Jun 2014 13:38:21 -0700 Subject: [PATCH 07/19] 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 fa17742349..d0097f3bee 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 72b5ab04c7..0854a92a88 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 55afd8c6b9..be4d8f1ee0 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; From 0e2ea45c93bde7e92a61736654aaff9cc7318e5b Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 08:46:14 -0700 Subject: [PATCH 08/19] ndisp no longer constant --- modules/cudastereo/src/cuda/stereocsbp.cu | 42 ++++++++++------------ modules/cudastereo/src/cuda/stereocsbp.hpp | 2 +- modules/cudastereo/src/stereocsbp.cpp | 2 +- 3 files changed, 21 insertions(+), 25 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index d0097f3bee..4c3bde337f 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -58,8 +58,6 @@ namespace cv { namespace cuda { namespace device /////////////////////// load constants //////////////////////// /////////////////////////////////////////////////////////////// - __constant__ int cndisp; - __constant__ float cmax_data_term; __constant__ float cdata_weight; __constant__ float cmax_disc_term; @@ -72,10 +70,8 @@ namespace cv { namespace cuda { namespace device __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) + void load_constants(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)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmax_data_term, &max_data_term, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cdata_weight, &data_weight, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cmax_disc_term, &max_disc_term, sizeof(float)) ); @@ -123,7 +119,7 @@ namespace cv { namespace cuda { namespace device }; template - __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane) + __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane, int ndisp) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -138,7 +134,7 @@ namespace cv { namespace cuda { namespace device { T minimum = device::numeric_limits::max(); int id = 0; - for(int d = 0; d < cndisp; d++) + for(int d = 0; d < ndisp; d++) { T cur = data_cost[d * cdisp_step1]; if(cur < minimum) @@ -157,7 +153,7 @@ namespace cv { namespace cuda { namespace device template - __global__ void get_first_k_initial_local(uchar *ctemp, T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane) + __global__ void get_first_k_initial_local(uchar *ctemp, T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane, int ndisp) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -174,7 +170,7 @@ namespace cv { namespace cuda { namespace device T cur = data_cost[1 * cdisp_step1]; T next = data_cost[2 * cdisp_step1]; - for (int d = 1; d < cndisp - 1 && nr_local_minimum < nr_plane; d++) + for (int d = 1; d < ndisp - 1 && nr_local_minimum < nr_plane; d++) { if (cur < prev && cur < next) { @@ -195,7 +191,7 @@ namespace cv { namespace cuda { namespace device T minimum = numeric_limits::max(); int id = 0; - for (int d = 0; d < cndisp; d++) + for (int d = 0; d < ndisp; d++) { cur = data_cost[d * cdisp_step1]; if (cur < minimum) @@ -213,7 +209,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, 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 ndisp) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -228,7 +224,7 @@ namespace cv { namespace cuda { namespace device T* data_cost = (T*)ctemp + y * cmsg_step + x; - for(int d = 0; d < cndisp; ++d) + for(int d = 0; d < ndisp; ++d) { float val = 0.0f; for(int yi = y0; yi < yt; yi++) @@ -253,7 +249,7 @@ namespace cv { namespace cuda { namespace device } template - __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) + __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 ndisp) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -261,7 +257,7 @@ namespace cv { namespace cuda { namespace device int tid = threadIdx.x; - if (d < cndisp) + if (d < ndisp) { int x0 = x_out << level; int y0 = y_out << level; @@ -301,7 +297,7 @@ namespace cv { namespace cuda { namespace device template - 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) + 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); @@ -311,9 +307,9 @@ namespace cv { namespace cuda { namespace device switch (channels) { - 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; + case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; + case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; + case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } @@ -330,9 +326,9 @@ namespace cv { namespace cuda { namespace device switch (channels) { - 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; + case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; + case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; + case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } @@ -368,9 +364,9 @@ namespace cv { namespace cuda { namespace device grid.y = divUp(h, threads.y); if (use_local_init_data_cost == true) - get_first_k_initial_local<<>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane); + get_first_k_initial_local<<>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp); else - get_first_k_initial_global<<>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane); + get_first_k_initial_global<<>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp); cudaSafeCall( cudaGetLastError() ); diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 0854a92a88..c9f3983256 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,7 +2,7 @@ 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); + void load_constants(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, size_t cimg_step, int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step, diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index be4d8f1ee0..946a14fb2b 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_); + load_constants(max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_); l[0].setTo(0, _stream); d[0].setTo(0, _stream); From 021b0cb4d53463b878b7b5a0bd1de3ccde5063db Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 09:28:26 -0700 Subject: [PATCH 09/19] Pass max_disc_term as kernel parameter. --- modules/cudastereo/src/cuda/stereocsbp.cu | 26 ++++++++++------------ modules/cudastereo/src/cuda/stereocsbp.hpp | 4 ++-- modules/cudastereo/src/stereocsbp.cpp | 6 ++--- 3 files changed, 17 insertions(+), 19 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 4c3bde337f..582aaa6440 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -60,7 +60,6 @@ namespace cv { namespace cuda { namespace device __constant__ float cmax_data_term; __constant__ float cdata_weight; - __constant__ float cmax_disc_term; __constant__ float cdisc_single_jump; __constant__ int cth; @@ -70,11 +69,10 @@ namespace cv { namespace cuda { namespace device __constant__ size_t cdisp_step2; - void load_constants(float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th) + void load_constants(float max_data_term, float data_weight, float disc_single_jump, int min_disp_th) { cudaSafeCall( cudaMemcpyToSymbol(cmax_data_term, &max_data_term, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cdata_weight, &data_weight, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmax_disc_term, &max_disc_term, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); @@ -688,7 +686,7 @@ namespace cv { namespace cuda { namespace device template __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3, - const T* dst_disp, const T* src_disp, int nr_plane, volatile T* temp) + const T* dst_disp, const T* src_disp, int nr_plane, int max_disc_term, volatile T* temp) { T minimum = numeric_limits::max(); @@ -706,7 +704,7 @@ namespace cv { namespace cuda { namespace device float sum = 0; for(int d = 0; d < nr_plane; d++) { - float cost_min = minimum + cmax_disc_term; + float cost_min = minimum + max_disc_term; T src_disp_reg = src_disp[d * cdisp_step1]; for(int d2 = 0; d2 < nr_plane; d2++) @@ -722,7 +720,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i) + __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i, int max_disc_term) { int y = blockIdx.y * blockDim.y + threadIdx.y; int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + i) & 1); @@ -740,17 +738,17 @@ namespace cv { namespace cuda { namespace device T* temp = (T*)ctemp + y * cmsg_step + x; - message_per_pixel(data, u, r - 1, u + cmsg_step, l + 1, disp, disp - cmsg_step, nr_plane, temp); - message_per_pixel(data, d, d - cmsg_step, r - 1, l + 1, disp, disp + cmsg_step, nr_plane, temp); - message_per_pixel(data, l, u + cmsg_step, d - cmsg_step, l + 1, disp, disp - 1, nr_plane, temp); - message_per_pixel(data, r, u + cmsg_step, d - cmsg_step, r - 1, disp, disp + 1, nr_plane, temp); + message_per_pixel(data, u, r - 1, u + cmsg_step, l + 1, disp, disp - cmsg_step, nr_plane, max_disc_term, temp); + message_per_pixel(data, d, d - cmsg_step, r - 1, l + 1, disp, disp + cmsg_step, nr_plane, max_disc_term, temp); + message_per_pixel(data, l, u + cmsg_step, d - cmsg_step, l + 1, disp, disp - 1, nr_plane, max_disc_term, temp); + message_per_pixel(data, r, u + cmsg_step, d - cmsg_step, r - 1, disp, disp + 1, nr_plane, max_disc_term, temp); } } template void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, - const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream) + const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream) { size_t disp_step = msg_step * h; cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); @@ -764,7 +762,7 @@ namespace cv { namespace cuda { namespace device for(int t = 0; t < iters; ++t) { - compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1); + compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1, max_disc_term); cudaSafeCall( cudaGetLastError() ); } if (stream == 0) @@ -772,10 +770,10 @@ namespace cv { namespace cuda { namespace device }; template void calc_all_iterations(uchar *ctemp, short* u, short* d, short* l, short* r, const short* data_cost_selected, const short* selected_disp_pyr_cur, size_t msg_step, - int h, int w, int nr_plane, int iters, cudaStream_t stream); + int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); template void calc_all_iterations(uchar *ctemp, float* u, float* d, float* l, float* r, const float* data_cost_selected, const float* selected_disp_pyr_cur, size_t msg_step, - int h, int w, int nr_plane, int iters, cudaStream_t stream); + int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); /////////////////////////////////////////////////////////////// diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index c9f3983256..80224754b9 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,7 +2,7 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - void load_constants(float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th); + void load_constants(float max_data_term, float data_weight, float disc_single_jump, int min_disp_th); template 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, @@ -21,7 +21,7 @@ namespace cv { namespace cuda { namespace device template void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, - const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream); + const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); template void compute_disp(const T* u, const T* d, const T* l, const T* r, const T* data_cost_selected, const T* disp_selected, size_t msg_step, diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 946a14fb2b..12d673153d 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -222,7 +222,7 @@ namespace //////////////////////////////////////////////////////////////////////////// // Compute - load_constants(max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_); + load_constants(max_data_term_, data_weight_, disc_single_jump_, min_disp_th_); l[0].setTo(0, _stream); d[0].setTo(0, _stream); @@ -267,7 +267,7 @@ namespace calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, - rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, stream); + rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, stream); } } else @@ -298,7 +298,7 @@ namespace calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, - rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, stream); + rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, stream); } } From 3ab117df04d361c77a56d4a5400404f264256ff4 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 09:51:32 -0700 Subject: [PATCH 10/19] Change struct with single static function to function. --- modules/cudastereo/src/cuda/stereocsbp.cu | 53 ++++++++++------------- 1 file changed, 22 insertions(+), 31 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 582aaa6440..3ef665b05c 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -82,39 +82,30 @@ namespace cv { namespace cuda { namespace device /////////////////////// init data cost //////////////////////// /////////////////////////////////////////////////////////////// - template struct DataCostPerPixel; - template <> struct DataCostPerPixel<1> + template float __device__ pixeldiff(const uchar* left, const uchar* right); + template<> float __device__ __forceinline__ pixeldiff<1>(const uchar* left, const uchar* right) { - static __device__ __forceinline__ float compute(const uchar* left, const uchar* right) - { - return cdata_weight * fmin( ::abs((int)*left - *right), cmax_data_term); - } - }; - template <> struct DataCostPerPixel<3> + return cdata_weight * fmin( ::abs((int)*left - *right), cmax_data_term); + } + template <> float __device__ __forceinline__ pixeldiff<3>(const uchar* left, const uchar* right) { - static __device__ __forceinline__ float compute(const uchar* left, const uchar* right) - { - float tb = 0.114f * ::abs((int)left[0] - right[0]); - float tg = 0.587f * ::abs((int)left[1] - right[1]); - float tr = 0.299f * ::abs((int)left[2] - right[2]); + float tb = 0.114f * ::abs((int)left[0] - right[0]); + float tg = 0.587f * ::abs((int)left[1] - right[1]); + float tr = 0.299f * ::abs((int)left[2] - right[2]); - return cdata_weight * fmin(tr + tg + tb, cmax_data_term); - } - }; - template <> struct DataCostPerPixel<4> + return cdata_weight * fmin(tr + tg + tb, cmax_data_term); + } + template <> float __device__ __forceinline__ pixeldiff<4>(const uchar* left, const uchar* right) { - static __device__ __forceinline__ float compute(const uchar* left, const uchar* right) - { - uchar4 l = *((const uchar4*)left); - uchar4 r = *((const uchar4*)right); + uchar4 l = *((const uchar4*)left); + uchar4 r = *((const uchar4*)right); - float tb = 0.114f * ::abs((int)l.x - r.x); - float tg = 0.587f * ::abs((int)l.y - r.y); - float tr = 0.299f * ::abs((int)l.z - r.z); + float tb = 0.114f * ::abs((int)l.x - r.x); + float tg = 0.587f * ::abs((int)l.y - r.y); + float tr = 0.299f * ::abs((int)l.z - r.z); - return cdata_weight * fmin(tr + tg + tb, cmax_data_term); - } - }; + return cdata_weight * fmin(tr + tg + tb, cmax_data_term); + } template __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane, int ndisp) @@ -237,7 +228,7 @@ namespace cv { namespace cuda { namespace device const uchar* lle = cleft + yi * cimg_step + xi * channels; const uchar* lri = cright + yi * cimg_step + xr * channels; - val += DataCostPerPixel::compute(lle, lri); + val += pixeldiff(lle, lri); } } } @@ -274,7 +265,7 @@ namespace cv { namespace cuda { namespace device for(int y = 0; y < len; ++y) { - val += DataCostPerPixel::compute(lle, lri); + val += pixeldiff(lle, lri); lle += cimg_step; lri += cimg_step; @@ -416,7 +407,7 @@ namespace cv { namespace cuda { namespace device const uchar* left_x = cleft + yi * cimg_step + xi * channels; const uchar* right_x = cright + yi * cimg_step + xr * channels; - val += DataCostPerPixel::compute(left_x, right_x); + val += pixeldiff(left_x, right_x); } } } @@ -458,7 +449,7 @@ namespace cv { namespace cuda { namespace device for(int y = 0; y < len; ++y) { - val += DataCostPerPixel::compute(lle, lri); + val += pixeldiff(lle, lri); lle += cimg_step; lri += cimg_step; From eed5cbc5db780e2c8ee5293f662cb663742c4e3e Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 13:22:13 -0700 Subject: [PATCH 11/19] More constant removal. --- modules/cudastereo/src/cuda/stereocsbp.cu | 96 +++++++++++----------- modules/cudastereo/src/cuda/stereocsbp.hpp | 6 +- modules/cudastereo/src/stereocsbp.cpp | 10 +-- 3 files changed, 54 insertions(+), 58 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 3ef665b05c..792df0c99c 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -58,8 +58,6 @@ namespace cv { namespace cuda { namespace device /////////////////////// load constants //////////////////////// /////////////////////////////////////////////////////////////// - __constant__ float cmax_data_term; - __constant__ float cdata_weight; __constant__ float cdisc_single_jump; __constant__ int cth; @@ -69,10 +67,8 @@ namespace cv { namespace cuda { namespace device __constant__ size_t cdisp_step2; - void load_constants(float max_data_term, float data_weight, float disc_single_jump, int min_disp_th) + void load_constants(float disc_single_jump, int min_disp_th) { - cudaSafeCall( cudaMemcpyToSymbol(cmax_data_term, &max_data_term, sizeof(float)) ); - cudaSafeCall( cudaMemcpyToSymbol(cdata_weight, &data_weight, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); @@ -82,20 +78,20 @@ namespace cv { namespace cuda { namespace device /////////////////////// init data cost //////////////////////// /////////////////////////////////////////////////////////////// - template float __device__ pixeldiff(const uchar* left, const uchar* right); - template<> float __device__ __forceinline__ pixeldiff<1>(const uchar* left, const uchar* right) + template static float __device__ pixeldiff(const uchar* left, const uchar* right, float max_data_term); + template<> __device__ __forceinline__ static float pixeldiff<1>(const uchar* left, const uchar* right, float max_data_term) { - return cdata_weight * fmin( ::abs((int)*left - *right), cmax_data_term); + return fmin( ::abs((int)*left - *right), max_data_term); } - template <> float __device__ __forceinline__ pixeldiff<3>(const uchar* left, const uchar* right) + template<> __device__ __forceinline__ static float pixeldiff<3>(const uchar* left, const uchar* right, float max_data_term) { float tb = 0.114f * ::abs((int)left[0] - right[0]); float tg = 0.587f * ::abs((int)left[1] - right[1]); float tr = 0.299f * ::abs((int)left[2] - right[2]); - return cdata_weight * fmin(tr + tg + tb, cmax_data_term); + return fmin(tr + tg + tb, max_data_term); } - template <> float __device__ __forceinline__ pixeldiff<4>(const uchar* left, const uchar* right) + template<> __device__ __forceinline__ static float pixeldiff<4>(const uchar* left, const uchar* right, float max_data_term) { uchar4 l = *((const uchar4*)left); uchar4 r = *((const uchar4*)right); @@ -104,7 +100,7 @@ namespace cv { namespace cuda { namespace device float tg = 0.587f * ::abs((int)l.y - r.y); float tr = 0.299f * ::abs((int)l.z - r.z); - return cdata_weight * fmin(tr + tg + tb, cmax_data_term); + return fmin(tr + tg + tb, max_data_term); } template @@ -198,7 +194,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level, int ndisp) + __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level, int ndisp, float data_weight, float max_data_term) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -222,13 +218,13 @@ namespace cv { namespace cuda { namespace device { int xr = xi - d; if(d < cth || xr < 0) - val += cdata_weight * cmax_data_term; + val += data_weight * max_data_term; else { const uchar* lle = cleft + yi * cimg_step + xi * channels; const uchar* lri = cright + yi * cimg_step + xr * channels; - val += pixeldiff(lle, lri); + val += data_weight * pixeldiff(lle, lri, max_data_term); } } } @@ -238,7 +234,7 @@ namespace cv { namespace cuda { namespace device } template - __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 ndisp) + __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 ndisp, float data_weight, float max_data_term) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -257,7 +253,7 @@ namespace cv { namespace cuda { namespace device if (x0 + tid < cols) { if (x0 + tid - d < 0 || d < cth) - val = cdata_weight * cmax_data_term * len; + val = data_weight * max_data_term * len; else { const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid ); @@ -265,7 +261,7 @@ namespace cv { namespace cuda { namespace device for(int y = 0; y < len; ++y) { - val += pixeldiff(lle, lri); + val += data_weight * pixeldiff(lle, lri, max_data_term); lle += cimg_step; lri += cimg_step; @@ -286,7 +282,7 @@ namespace cv { namespace cuda { namespace device template - 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) + 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, float data_weight, float max_data_term, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -296,15 +292,15 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; - case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; - case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break; + case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break; + case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break; + case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); 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, size_t cimg_step, 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, float data_weight, float max_data_term, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -315,19 +311,19 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; - case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; - case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break; + case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; + case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; + case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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) + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, bool use_local_init_data_cost, 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); + 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, float data_weight, float max_data_term, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -340,7 +336,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, cimg_step, 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, data_weight, max_data_term, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -364,17 +360,17 @@ namespace cv { namespace cuda { namespace device } 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); + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); 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); + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); /////////////////////////////////////////////////////////////// ////////////////////// compute data cost ////////////////////// /////////////////////////////////////////////////////////////// template - __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) + __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, float data_weight, float max_data_term) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -401,13 +397,13 @@ namespace cv { namespace cuda { namespace device int xr = xi - sel_disp; if (xr < 0 || sel_disp < cth) - val += cdata_weight * cmax_data_term; + val += data_weight * max_data_term; else { const uchar* left_x = cleft + yi * cimg_step + xi * channels; const uchar* right_x = cright + yi * cimg_step + xr * channels; - val += pixeldiff(left_x, right_x); + val += data_weight * pixeldiff(left_x, right_x, max_data_term); } } } @@ -417,7 +413,7 @@ namespace cv { namespace cuda { namespace device } template - __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) + __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, float data_weight, float max_data_term) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -441,7 +437,7 @@ namespace cv { namespace cuda { namespace device if (x0 + tid < cols) { if (x0 + tid - sel_disp < 0 || sel_disp < cth) - val = cdata_weight * cmax_data_term * len; + val = data_weight * max_data_term * len; else { const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid ); @@ -449,7 +445,7 @@ namespace cv { namespace cuda { namespace device for(int y = 0; y < len; ++y) { - val += pixeldiff(lle, lri); + val += data_weight * pixeldiff(lle, lri, max_data_term); lle += cimg_step; lri += cimg_step; @@ -468,7 +464,7 @@ namespace cv { namespace cuda { namespace device template 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) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -478,16 +474,16 @@ namespace cv { namespace cuda { namespace device switch(channels) { - 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; + case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; + case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; + case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -498,19 +494,19 @@ namespace cv { namespace cuda { namespace device switch (channels) { - 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; + case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; + case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; + case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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) + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream) { 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); + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream); static const ComputeDataCostCaller callers[] = { @@ -525,7 +521,7 @@ 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, cimg_step, 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, data_weight, max_data_term, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -533,10 +529,10 @@ namespace cv { namespace cuda { namespace device } 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); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream); 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); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream); /////////////////////////////////////////////////////////////// diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 80224754b9..04337bbd1d 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,15 +2,15 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - void load_constants(float max_data_term, float data_weight, float disc_single_jump, int min_disp_th); + void load_constants(float disc_single_jump, int min_disp_th); template 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); + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); template 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); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, cudaStream_t stream); template void init_message(uchar *ctemp, T* u_new, T* d_new, T* l_new, T* r_new, diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 12d673153d..14d6d2589f 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -222,7 +222,7 @@ namespace //////////////////////////////////////////////////////////////////////////// // Compute - load_constants(max_data_term_, data_weight_, disc_single_jump_, min_disp_th_); + load_constants(disc_single_jump_, min_disp_th_); l[0].setTo(0, _stream); d[0].setTo(0, _stream); @@ -246,12 +246,12 @@ namespace if (i == levels_ - 1) { 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); + elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), data_weight_, max_data_term_, use_local_init_data_cost_, stream); } else { 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); + left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), data_weight_, max_data_term_, stream); int new_idx = (cur_idx + 1) & 1; @@ -277,12 +277,12 @@ namespace if (i == levels_ - 1) { 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); + elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), data_weight_, max_data_term_, use_local_init_data_cost_, stream); } else { 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); + left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), data_weight_, max_data_term_, stream); int new_idx = (cur_idx + 1) & 1; From 9bc71f4cb61ac82ce166e4638b6cc0b8d5d5c5ae Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 14:20:46 -0700 Subject: [PATCH 12/19] Deconstify minimum disparity. --- modules/cudastereo/src/cuda/stereocsbp.cu | 79 +++++++++++----------- modules/cudastereo/src/cuda/stereocsbp.hpp | 7 +- modules/cudastereo/src/stereocsbp.cpp | 10 +-- 3 files changed, 47 insertions(+), 49 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 792df0c99c..fc6f4f3daf 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -60,18 +60,14 @@ namespace cv { namespace cuda { namespace device __constant__ float cdisc_single_jump; - __constant__ int cth; - __constant__ size_t cmsg_step; __constant__ size_t cdisp_step1; __constant__ size_t cdisp_step2; - void load_constants(float disc_single_jump, int min_disp_th) + void load_constants(float disc_single_jump) { cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); - - cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); } /////////////////////////////////////////////////////////////// @@ -194,7 +190,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level, int ndisp, float data_weight, float max_data_term) + __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level, int ndisp, float data_weight, float max_data_term, int min_disp) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -217,7 +213,7 @@ namespace cv { namespace cuda { namespace device for(int xi = x0; xi < xt; xi++) { int xr = xi - d; - if(d < cth || xr < 0) + if(d < min_disp || xr < 0) val += data_weight * max_data_term; else { @@ -234,7 +230,7 @@ namespace cv { namespace cuda { namespace device } template - __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 ndisp, float data_weight, float max_data_term) + __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 ndisp, float data_weight, float max_data_term, int min_disp) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -252,7 +248,7 @@ namespace cv { namespace cuda { namespace device float val = 0.0f; if (x0 + tid < cols) { - if (x0 + tid - d < 0 || d < cth) + if (x0 + tid - d < 0 || d < min_disp) val = data_weight * max_data_term * len; else { @@ -282,7 +278,7 @@ namespace cv { namespace cuda { namespace device template - 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, float data_weight, float max_data_term, 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -292,15 +288,15 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break; - case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break; - case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break; + case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); break; + case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); break; + case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); 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, size_t cimg_step, int rows, int cols, int h, int w, int level, int ndisp, int channels, float data_weight, float max_data_term, 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -311,19 +307,19 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; - case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; - case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break; + case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; + case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; + case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream) + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, 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, float data_weight, float max_data_term, 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -336,7 +332,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, cimg_step, rows, cols, h, w, level, ndisp, channels, data_weight, max_data_term, stream); + init_data_cost_callers[level](cleft, cright, ctemp, cimg_step, rows, cols, h, w, level, ndisp, channels, data_weight, max_data_term, min_disp, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -359,18 +355,18 @@ namespace cv { namespace cuda { namespace device cudaSafeCall( cudaDeviceSynchronize() ); } - 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, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); + 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, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, cudaStream_t stream); - 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, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); + 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, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, cudaStream_t stream); /////////////////////////////////////////////////////////////// ////////////////////// compute data cost ////////////////////// /////////////////////////////////////////////////////////////// template - __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, float data_weight, float max_data_term) + __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, float data_weight, float max_data_term, int min_disp) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -396,7 +392,7 @@ namespace cv { namespace cuda { namespace device int sel_disp = selected_disparity[d * cdisp_step2]; int xr = xi - sel_disp; - if (xr < 0 || sel_disp < cth) + if (xr < 0 || sel_disp < min_disp) val += data_weight * max_data_term; else { @@ -413,7 +409,7 @@ namespace cv { namespace cuda { namespace device } template - __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, float data_weight, float max_data_term) + __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, float data_weight, float max_data_term, int min_disp) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -436,7 +432,7 @@ namespace cv { namespace cuda { namespace device float val = 0.0f; if (x0 + tid < cols) { - if (x0 + tid - sel_disp < 0 || sel_disp < cth) + if (x0 + tid - sel_disp < 0 || sel_disp < min_disp) val = data_weight * max_data_term * len; else { @@ -464,7 +460,7 @@ namespace cv { namespace cuda { namespace device template 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, float data_weight, float max_data_term, cudaStream_t stream) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -474,16 +470,16 @@ namespace cv { namespace cuda { namespace device switch(channels) { - case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; - case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; - case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); break; + case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; + case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; + case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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, float data_weight, float max_data_term, cudaStream_t stream) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -494,19 +490,20 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; - case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; - case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term); break; + case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; + case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; + case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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, float data_weight, float max_data_term, cudaStream_t stream) + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, + int min_disp, cudaStream_t stream) { 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, float data_weight, float max_data_term, cudaStream_t stream); + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, cudaStream_t stream); static const ComputeDataCostCaller callers[] = { @@ -521,7 +518,7 @@ 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, cimg_step, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, data_weight, max_data_term, stream); + callers[level](cleft, cright, cimg_step, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, data_weight, max_data_term, min_disp, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -529,10 +526,10 @@ namespace cv { namespace cuda { namespace device } 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, float data_weight, float max_data_term, cudaStream_t stream); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, cudaStream_t stream); 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, float data_weight, float max_data_term, cudaStream_t stream); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, cudaStream_t stream); /////////////////////////////////////////////////////////////// diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 04337bbd1d..95c5a47356 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,15 +2,16 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - void load_constants(float disc_single_jump, int min_disp_th); + void load_constants(float disc_single_jump); template 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, float data_weight, float max_data_term, bool use_local_init_data_cost, cudaStream_t stream); + int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, cudaStream_t stream); template 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, float data_weight, float max_data_term, cudaStream_t stream); + int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, float data_weight, float max_data_term, + int min_disp, cudaStream_t stream); template void init_message(uchar *ctemp, T* u_new, T* d_new, T* l_new, T* r_new, diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 14d6d2589f..2515061e13 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -222,7 +222,7 @@ namespace //////////////////////////////////////////////////////////////////////////// // Compute - load_constants(disc_single_jump_, min_disp_th_); + load_constants(disc_single_jump_); l[0].setTo(0, _stream); d[0].setTo(0, _stream); @@ -246,12 +246,12 @@ namespace if (i == levels_ - 1) { 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(), data_weight_, max_data_term_, use_local_init_data_cost_, stream); + elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), data_weight_, max_data_term_, min_disp_th_, use_local_init_data_cost_, stream); } else { 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(), data_weight_, max_data_term_, stream); + left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), data_weight_, max_data_term_, min_disp_th_, stream); int new_idx = (cur_idx + 1) & 1; @@ -277,12 +277,12 @@ namespace if (i == levels_ - 1) { 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(), data_weight_, max_data_term_, use_local_init_data_cost_, stream); + elem_step, rows_pyr[i], cols_pyr[i], i, nr_plane_pyr[i], ndisp_, left.channels(), data_weight_, max_data_term_, min_disp_th_, use_local_init_data_cost_, stream); } else { 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(), data_weight_, max_data_term_, stream); + left.rows, left.cols, rows_pyr[i], cols_pyr[i], rows_pyr[i+1], i, nr_plane_pyr[i+1], left.channels(), data_weight_, max_data_term_, min_disp_th_, stream); int new_idx = (cur_idx + 1) & 1; From 6d86d63ac59a9e581c0680c8a9580f75d9d98d4f Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Tue, 1 Jul 2014 10:44:52 -0700 Subject: [PATCH 13/19] Last of csbp load_constants() gone. --- modules/cudastereo/src/cuda/stereocsbp.cu | 30 ++++++++-------------- modules/cudastereo/src/cuda/stereocsbp.hpp | 4 +-- modules/cudastereo/src/stereocsbp.cpp | 6 ++--- 3 files changed, 14 insertions(+), 26 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index fc6f4f3daf..04f6caced6 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -58,18 +58,10 @@ namespace cv { namespace cuda { namespace device /////////////////////// load constants //////////////////////// /////////////////////////////////////////////////////////////// - __constant__ float cdisc_single_jump; - __constant__ size_t cmsg_step; __constant__ size_t cdisp_step1; __constant__ size_t cdisp_step2; - - void load_constants(float disc_single_jump) - { - cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); - } - /////////////////////////////////////////////////////////////// /////////////////////// init data cost //////////////////////// /////////////////////////////////////////////////////////////// @@ -670,7 +662,7 @@ namespace cv { namespace cuda { namespace device template __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3, - const T* dst_disp, const T* src_disp, int nr_plane, int max_disc_term, volatile T* temp) + const T* dst_disp, const T* src_disp, int nr_plane, int max_disc_term, float disc_single_jump, volatile T* temp) { T minimum = numeric_limits::max(); @@ -692,7 +684,7 @@ namespace cv { namespace cuda { namespace device T src_disp_reg = src_disp[d * cdisp_step1]; for(int d2 = 0; d2 < nr_plane; d2++) - cost_min = fmin(cost_min, msg_dst[d2 * cdisp_step1] + cdisc_single_jump * ::abs(dst_disp[d2 * cdisp_step1] - src_disp_reg)); + cost_min = fmin(cost_min, msg_dst[d2 * cdisp_step1] + disc_single_jump * ::abs(dst_disp[d2 * cdisp_step1] - src_disp_reg)); temp[d * cdisp_step1] = saturate_cast(cost_min); sum += cost_min; @@ -704,7 +696,7 @@ namespace cv { namespace cuda { namespace device } template - __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i, int max_disc_term) + __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i, int max_disc_term, float disc_single_jump) { int y = blockIdx.y * blockDim.y + threadIdx.y; int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + i) & 1); @@ -722,17 +714,17 @@ namespace cv { namespace cuda { namespace device T* temp = (T*)ctemp + y * cmsg_step + x; - message_per_pixel(data, u, r - 1, u + cmsg_step, l + 1, disp, disp - cmsg_step, nr_plane, max_disc_term, temp); - message_per_pixel(data, d, d - cmsg_step, r - 1, l + 1, disp, disp + cmsg_step, nr_plane, max_disc_term, temp); - message_per_pixel(data, l, u + cmsg_step, d - cmsg_step, l + 1, disp, disp - 1, nr_plane, max_disc_term, temp); - message_per_pixel(data, r, u + cmsg_step, d - cmsg_step, r - 1, disp, disp + 1, nr_plane, max_disc_term, temp); + message_per_pixel(data, u, r - 1, u + cmsg_step, l + 1, disp, disp - cmsg_step, nr_plane, max_disc_term, disc_single_jump, temp); + message_per_pixel(data, d, d - cmsg_step, r - 1, l + 1, disp, disp + cmsg_step, nr_plane, max_disc_term, disc_single_jump, temp); + message_per_pixel(data, l, u + cmsg_step, d - cmsg_step, l + 1, disp, disp - 1, nr_plane, max_disc_term, disc_single_jump, temp); + message_per_pixel(data, r, u + cmsg_step, d - cmsg_step, r - 1, disp, disp + 1, nr_plane, max_disc_term, disc_single_jump, temp); } } template void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, - const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream) + const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream) { size_t disp_step = msg_step * h; cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); @@ -746,7 +738,7 @@ namespace cv { namespace cuda { namespace device for(int t = 0; t < iters; ++t) { - compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1, max_disc_term); + compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1, max_disc_term, disc_single_jump); cudaSafeCall( cudaGetLastError() ); } if (stream == 0) @@ -754,10 +746,10 @@ namespace cv { namespace cuda { namespace device }; template void calc_all_iterations(uchar *ctemp, short* u, short* d, short* l, short* r, const short* data_cost_selected, const short* selected_disp_pyr_cur, size_t msg_step, - int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); + int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream); template void calc_all_iterations(uchar *ctemp, float* u, float* d, float* l, float* r, const float* data_cost_selected, const float* selected_disp_pyr_cur, size_t msg_step, - int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); + int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream); /////////////////////////////////////////////////////////////// diff --git a/modules/cudastereo/src/cuda/stereocsbp.hpp b/modules/cudastereo/src/cuda/stereocsbp.hpp index 95c5a47356..305497292d 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.hpp +++ b/modules/cudastereo/src/cuda/stereocsbp.hpp @@ -2,8 +2,6 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - void load_constants(float disc_single_jump); - template 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, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, cudaStream_t stream); @@ -22,7 +20,7 @@ namespace cv { namespace cuda { namespace device template void calc_all_iterations(uchar *ctemp, T* u, T* d, T* l, T* r, const T* data_cost_selected, - const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, cudaStream_t stream); + const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream); template void compute_disp(const T* u, const T* d, const T* l, const T* r, const T* data_cost_selected, const T* disp_selected, size_t msg_step, diff --git a/modules/cudastereo/src/stereocsbp.cpp b/modules/cudastereo/src/stereocsbp.cpp index 2515061e13..ded5fa20e1 100644 --- a/modules/cudastereo/src/stereocsbp.cpp +++ b/modules/cudastereo/src/stereocsbp.cpp @@ -222,8 +222,6 @@ namespace //////////////////////////////////////////////////////////////////////////// // Compute - load_constants(disc_single_jump_); - l[0].setTo(0, _stream); d[0].setTo(0, _stream); r[0].setTo(0, _stream); @@ -267,7 +265,7 @@ namespace calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, - rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, stream); + rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, disc_single_jump_, stream); } } else @@ -298,7 +296,7 @@ namespace calc_all_iterations(temp_.ptr(), u[cur_idx].ptr(), d[cur_idx].ptr(), l[cur_idx].ptr(), r[cur_idx].ptr(), data_cost_selected.ptr(), disp_selected_pyr[cur_idx].ptr(), elem_step, - rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, stream); + rows_pyr[i], cols_pyr[i], nr_plane_pyr[i], iters_, max_disc_term_, disc_single_jump_, stream); } } From b792419cde59876969de2f803c1599920cf2a0f2 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 7 Jul 2014 08:12:28 -0700 Subject: [PATCH 14/19] Remove compute_disp()'s use of constant memory. --- modules/cudastereo/src/cuda/stereocsbp.cu | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 04f6caced6..79456a01fa 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -760,26 +760,26 @@ namespace cv { namespace cuda { namespace device template __global__ void compute_disp(const T* u_, const T* d_, const T* l_, const T* r_, const T* data_cost_selected, const T* disp_selected_pyr, - PtrStepSz disp, int nr_plane) + PtrStepSz disp, int nr_plane, size_t msg_step, size_t disp_step) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y > 0 && y < disp.rows - 1 && x > 0 && x < disp.cols - 1) { - const T* data = data_cost_selected + y * cmsg_step + x; - const T* disp_selected = disp_selected_pyr + y * cmsg_step + x; + const T* data = data_cost_selected + y * msg_step + x; + const T* disp_selected = disp_selected_pyr + y * msg_step + x; - const T* u = u_ + (y+1) * cmsg_step + (x+0); - const T* d = d_ + (y-1) * cmsg_step + (x+0); - const T* l = l_ + (y+0) * cmsg_step + (x+1); - const T* r = r_ + (y+0) * cmsg_step + (x-1); + const T* u = u_ + (y+1) * msg_step + (x+0); + const T* d = d_ + (y-1) * msg_step + (x+0); + const T* l = l_ + (y+0) * msg_step + (x+1); + const T* r = r_ + (y+0) * msg_step + (x-1); int best = 0; T best_val = numeric_limits::max(); for (int i = 0; i < nr_plane; ++i) { - int idx = i * cdisp_step1; + int idx = i * disp_step; T val = data[idx]+ u[idx] + d[idx] + l[idx] + r[idx]; if (val < best_val) @@ -797,8 +797,6 @@ namespace cv { namespace cuda { namespace device const PtrStepSz& disp, int nr_plane, cudaStream_t stream) { size_t disp_step = disp.rows * msg_step; - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -806,7 +804,7 @@ namespace cv { namespace cuda { namespace device grid.x = divUp(disp.cols, threads.x); grid.y = divUp(disp.rows, threads.y); - compute_disp<<>>(u, d, l, r, data_cost_selected, disp_selected, disp, nr_plane); + compute_disp<<>>(u, d, l, r, data_cost_selected, disp_selected, disp, nr_plane, msg_step, disp_step); cudaSafeCall( cudaGetLastError() ); if (stream == 0) From 9b8002cd43c8735e6e2eb1ecab1930bad6325681 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 7 Jul 2014 08:23:40 -0700 Subject: [PATCH 15/19] remove use of constant memory in calc_all_iterations/compute_message/message_per_pixel --- modules/cudastereo/src/cuda/stereocsbp.cu | 41 +++++++++++------------ 1 file changed, 20 insertions(+), 21 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 79456a01fa..8c38d91ae9 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -662,13 +662,14 @@ namespace cv { namespace cuda { namespace device template __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3, - const T* dst_disp, const T* src_disp, int nr_plane, int max_disc_term, float disc_single_jump, volatile T* temp) + const T* dst_disp, const T* src_disp, int nr_plane, int max_disc_term, float disc_single_jump, volatile T* temp, + size_t disp_step) { T minimum = numeric_limits::max(); for(int d = 0; d < nr_plane; d++) { - int idx = d * cdisp_step1; + int idx = d * disp_step; T val = data[idx] + msg1[idx] + msg2[idx] + msg3[idx]; if(val < minimum) @@ -681,43 +682,43 @@ namespace cv { namespace cuda { namespace device for(int d = 0; d < nr_plane; d++) { float cost_min = minimum + max_disc_term; - T src_disp_reg = src_disp[d * cdisp_step1]; + T src_disp_reg = src_disp[d * disp_step]; for(int d2 = 0; d2 < nr_plane; d2++) - cost_min = fmin(cost_min, msg_dst[d2 * cdisp_step1] + disc_single_jump * ::abs(dst_disp[d2 * cdisp_step1] - src_disp_reg)); + cost_min = fmin(cost_min, msg_dst[d2 * disp_step] + disc_single_jump * ::abs(dst_disp[d2 * disp_step] - src_disp_reg)); - temp[d * cdisp_step1] = saturate_cast(cost_min); + temp[d * disp_step] = saturate_cast(cost_min); sum += cost_min; } sum /= nr_plane; for(int d = 0; d < nr_plane; d++) - msg_dst[d * cdisp_step1] = saturate_cast(temp[d * cdisp_step1] - sum); + msg_dst[d * disp_step] = saturate_cast(temp[d * disp_step] - sum); } template - __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i, int max_disc_term, float disc_single_jump) + __global__ void compute_message(uchar *ctemp, T* u_, T* d_, T* l_, T* r_, const T* data_cost_selected, const T* selected_disp_pyr_cur, int h, int w, int nr_plane, int i, int max_disc_term, float disc_single_jump, size_t msg_step, size_t disp_step) { int y = blockIdx.y * blockDim.y + threadIdx.y; int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + i) & 1); if (y > 0 && y < h - 1 && x > 0 && x < w - 1) { - const T* data = data_cost_selected + y * cmsg_step + x; + const T* data = data_cost_selected + y * msg_step + x; - T* u = u_ + y * cmsg_step + x; - T* d = d_ + y * cmsg_step + x; - T* l = l_ + y * cmsg_step + x; - T* r = r_ + y * cmsg_step + x; + T* u = u_ + y * msg_step + x; + T* d = d_ + y * msg_step + x; + T* l = l_ + y * msg_step + x; + T* r = r_ + y * msg_step + x; - const T* disp = selected_disp_pyr_cur + y * cmsg_step + x; + const T* disp = selected_disp_pyr_cur + y * msg_step + x; - T* temp = (T*)ctemp + y * cmsg_step + x; + T* temp = (T*)ctemp + y * msg_step + x; - message_per_pixel(data, u, r - 1, u + cmsg_step, l + 1, disp, disp - cmsg_step, nr_plane, max_disc_term, disc_single_jump, temp); - message_per_pixel(data, d, d - cmsg_step, r - 1, l + 1, disp, disp + cmsg_step, nr_plane, max_disc_term, disc_single_jump, temp); - message_per_pixel(data, l, u + cmsg_step, d - cmsg_step, l + 1, disp, disp - 1, nr_plane, max_disc_term, disc_single_jump, temp); - message_per_pixel(data, r, u + cmsg_step, d - cmsg_step, r - 1, disp, disp + 1, nr_plane, max_disc_term, disc_single_jump, temp); + message_per_pixel(data, u, r - 1, u + msg_step, l + 1, disp, disp - msg_step, nr_plane, max_disc_term, disc_single_jump, temp, disp_step); + message_per_pixel(data, d, d - msg_step, r - 1, l + 1, disp, disp + msg_step, nr_plane, max_disc_term, disc_single_jump, temp, disp_step); + message_per_pixel(data, l, u + msg_step, d - msg_step, l + 1, disp, disp - 1, nr_plane, max_disc_term, disc_single_jump, temp, disp_step); + message_per_pixel(data, r, u + msg_step, d - msg_step, r - 1, disp, disp + 1, nr_plane, max_disc_term, disc_single_jump, temp, disp_step); } } @@ -727,8 +728,6 @@ namespace cv { namespace cuda { namespace device const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, int max_disc_term, float disc_single_jump, cudaStream_t stream) { size_t disp_step = msg_step * h; - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -738,7 +737,7 @@ namespace cv { namespace cuda { namespace device for(int t = 0; t < iters; ++t) { - compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1, max_disc_term, disc_single_jump); + compute_message<<>>(ctemp, u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1, max_disc_term, disc_single_jump, msg_step, disp_step); cudaSafeCall( cudaGetLastError() ); } if (stream == 0) From 1ff270e41c4c69952361baf87ddf304ea0151a0e Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 7 Jul 2014 09:45:30 -0700 Subject: [PATCH 16/19] init_message no longer uses constant memory. --- modules/cudastereo/src/cuda/stereocsbp.cu | 68 +++++++++++------------ 1 file changed, 34 insertions(+), 34 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 8c38d91ae9..6ebdee8e9a 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -534,7 +534,7 @@ namespace cv { namespace cuda { namespace device const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, T* data_cost_selected, T* disparity_selected_new, T* data_cost_new, const T* data_cost_cur, const T* disparity_selected_cur, - int nr_plane, int nr_plane2) + int nr_plane, int nr_plane2, size_t disp_step1, size_t disp_step2) { for(int i = 0; i < nr_plane; i++) { @@ -550,15 +550,15 @@ namespace cv { namespace cuda { namespace device } } - data_cost_selected[i * cdisp_step1] = data_cost_cur[id * cdisp_step1]; - disparity_selected_new[i * cdisp_step1] = disparity_selected_cur[id * cdisp_step2]; + data_cost_selected[i * disp_step1] = data_cost_cur[id * disp_step1]; + disparity_selected_new[i * disp_step1] = disparity_selected_cur[id * disp_step2]; - u_new[i * cdisp_step1] = u_cur[id * cdisp_step2]; - d_new[i * cdisp_step1] = d_cur[id * cdisp_step2]; - l_new[i * cdisp_step1] = l_cur[id * cdisp_step2]; - r_new[i * cdisp_step1] = r_cur[id * cdisp_step2]; + u_new[i * disp_step1] = u_cur[id * disp_step2]; + d_new[i * disp_step1] = d_cur[id * disp_step2]; + l_new[i * disp_step1] = l_cur[id * disp_step2]; + r_new[i * disp_step1] = r_cur[id * disp_step2]; - data_cost_new[id * cdisp_step1] = numeric_limits::max(); + data_cost_new[id * disp_step1] = numeric_limits::max(); } } @@ -567,47 +567,49 @@ namespace cv { namespace cuda { namespace device const T* u_cur_, const T* d_cur_, const T* l_cur_, const T* r_cur_, T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, T* data_cost_selected_, const T* data_cost_, - int h, int w, int nr_plane, int h2, int w2, int nr_plane2) + int h, int w, int nr_plane, int h2, int w2, int nr_plane2, + size_t msg_step, size_t disp_step1, size_t disp_step2) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < h && x < w) { - const T* u_cur = u_cur_ + ::min(h2-1, y/2 + 1) * cmsg_step + x/2; - const T* d_cur = d_cur_ + ::max(0, y/2 - 1) * cmsg_step + x/2; - const T* l_cur = l_cur_ + (y/2) * cmsg_step + ::min(w2-1, x/2 + 1); - const T* r_cur = r_cur_ + (y/2) * cmsg_step + ::max(0, x/2 - 1); + const T* u_cur = u_cur_ + ::min(h2-1, y/2 + 1) * msg_step + x/2; + const T* d_cur = d_cur_ + ::max(0, y/2 - 1) * msg_step + x/2; + const T* l_cur = l_cur_ + (y/2) * msg_step + ::min(w2-1, x/2 + 1); + const T* r_cur = r_cur_ + (y/2) * msg_step + ::max(0, x/2 - 1); - T* data_cost_new = (T*)ctemp + y * cmsg_step + x; + T* data_cost_new = (T*)ctemp + y * msg_step + x; - const T* disparity_selected_cur = selected_disp_pyr_cur + y/2 * cmsg_step + x/2; - const T* data_cost = data_cost_ + y * cmsg_step + x; + const T* disparity_selected_cur = selected_disp_pyr_cur + y/2 * msg_step + x/2; + const T* data_cost = data_cost_ + y * msg_step + x; for(int d = 0; d < nr_plane2; d++) { - int idx2 = d * cdisp_step2; + int idx2 = d * disp_step2; - T val = data_cost[d * cdisp_step1] + u_cur[idx2] + d_cur[idx2] + l_cur[idx2] + r_cur[idx2]; - data_cost_new[d * cdisp_step1] = val; + T val = data_cost[d * disp_step1] + u_cur[idx2] + d_cur[idx2] + l_cur[idx2] + r_cur[idx2]; + data_cost_new[d * disp_step1] = val; } - T* data_cost_selected = data_cost_selected_ + y * cmsg_step + x; - T* disparity_selected_new = selected_disp_pyr_new + y * cmsg_step + x; + T* data_cost_selected = data_cost_selected_ + y * msg_step + x; + T* disparity_selected_new = selected_disp_pyr_new + y * msg_step + x; - T* u_new = u_new_ + y * cmsg_step + x; - T* d_new = d_new_ + y * cmsg_step + x; - T* l_new = l_new_ + y * cmsg_step + x; - T* r_new = r_new_ + y * cmsg_step + x; + T* u_new = u_new_ + y * msg_step + x; + T* d_new = d_new_ + y * msg_step + x; + T* l_new = l_new_ + y * msg_step + x; + T* r_new = r_new_ + y * msg_step + x; - u_cur = u_cur_ + y/2 * cmsg_step + x/2; - d_cur = d_cur_ + y/2 * cmsg_step + x/2; - l_cur = l_cur_ + y/2 * cmsg_step + x/2; - r_cur = r_cur_ + y/2 * cmsg_step + x/2; + u_cur = u_cur_ + y/2 * msg_step + x/2; + d_cur = d_cur_ + y/2 * msg_step + x/2; + l_cur = l_cur_ + y/2 * msg_step + x/2; + r_cur = r_cur_ + y/2 * msg_step + x/2; get_first_k_element_increase(u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur, data_cost_selected, disparity_selected_new, data_cost_new, - data_cost, disparity_selected_cur, nr_plane, nr_plane2); + data_cost, disparity_selected_cur, nr_plane, nr_plane2, + disp_step1, disp_step2); } } @@ -622,9 +624,6 @@ namespace cv { namespace cuda { namespace device size_t disp_step1 = msg_step * h; size_t disp_step2 = msg_step * h2; - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step1, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step2, &disp_step2, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -636,7 +635,8 @@ namespace cv { namespace cuda { namespace device u_cur, d_cur, l_cur, r_cur, selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, - h, w, nr_plane, h2, w2, nr_plane2); + h, w, nr_plane, h2, w2, nr_plane2, + msg_step, disp_step1, disp_step2); cudaSafeCall( cudaGetLastError() ); if (stream == 0) From 52516085d98e3ca19d2785143acfe992b78c29a6 Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 7 Jul 2014 10:12:30 -0700 Subject: [PATCH 17/19] remove constant memory from init_data_cost --- modules/cudastereo/src/cuda/stereocsbp.cu | 92 ++++++++++++----------- 1 file changed, 48 insertions(+), 44 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 6ebdee8e9a..974b503733 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -92,16 +92,17 @@ namespace cv { namespace cuda { namespace device } template - __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane, int ndisp) + __global__ void get_first_k_initial_global(uchar *ctemp, T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane, int ndisp, + size_t msg_step, size_t disp_step) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < h && x < w) { - T* selected_disparity = selected_disp_pyr + y * cmsg_step + x; - T* data_cost_selected = data_cost_selected_ + y * cmsg_step + x; - T* data_cost = (T*)ctemp + y * cmsg_step + x; + T* selected_disparity = selected_disp_pyr + y * msg_step + x; + T* data_cost_selected = data_cost_selected_ + y * msg_step + x; + T* data_cost = (T*)ctemp + y * msg_step + x; for(int i = 0; i < nr_plane; i++) { @@ -109,7 +110,7 @@ namespace cv { namespace cuda { namespace device int id = 0; for(int d = 0; d < ndisp; d++) { - T cur = data_cost[d * cdisp_step1]; + T cur = data_cost[d * disp_step]; if(cur < minimum) { minimum = cur; @@ -117,46 +118,47 @@ namespace cv { namespace cuda { namespace device } } - data_cost_selected[i * cdisp_step1] = minimum; - selected_disparity[i * cdisp_step1] = id; - data_cost [id * cdisp_step1] = numeric_limits::max(); + data_cost_selected[i * disp_step] = minimum; + selected_disparity[i * disp_step] = id; + data_cost [id * disp_step] = numeric_limits::max(); } } } template - __global__ void get_first_k_initial_local(uchar *ctemp, T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane, int ndisp) + __global__ void get_first_k_initial_local(uchar *ctemp, T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane, int ndisp, + size_t msg_step, size_t disp_step) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; if (y < h && x < w) { - T* selected_disparity = selected_disp_pyr + y * cmsg_step + x; - T* data_cost_selected = data_cost_selected_ + y * cmsg_step + x; - T* data_cost = (T*)ctemp + y * cmsg_step + x; + T* selected_disparity = selected_disp_pyr + y * msg_step + x; + T* data_cost_selected = data_cost_selected_ + y * msg_step + x; + T* data_cost = (T*)ctemp + y * msg_step + x; int nr_local_minimum = 0; - T prev = data_cost[0 * cdisp_step1]; - T cur = data_cost[1 * cdisp_step1]; - T next = data_cost[2 * cdisp_step1]; + T prev = data_cost[0 * disp_step]; + T cur = data_cost[1 * disp_step]; + T next = data_cost[2 * disp_step]; for (int d = 1; d < ndisp - 1 && nr_local_minimum < nr_plane; d++) { if (cur < prev && cur < next) { - data_cost_selected[nr_local_minimum * cdisp_step1] = cur; - selected_disparity[nr_local_minimum * cdisp_step1] = d; + data_cost_selected[nr_local_minimum * disp_step] = cur; + selected_disparity[nr_local_minimum * disp_step] = d; - data_cost[d * cdisp_step1] = numeric_limits::max(); + data_cost[d * disp_step] = numeric_limits::max(); nr_local_minimum++; } prev = cur; cur = next; - next = data_cost[(d + 1) * cdisp_step1]; + next = data_cost[(d + 1) * disp_step]; } for (int i = nr_local_minimum; i < nr_plane; i++) @@ -166,23 +168,25 @@ namespace cv { namespace cuda { namespace device for (int d = 0; d < ndisp; d++) { - cur = data_cost[d * cdisp_step1]; + cur = data_cost[d * disp_step]; if (cur < minimum) { minimum = cur; id = d; } } - data_cost_selected[i * cdisp_step1] = minimum; - selected_disparity[i * cdisp_step1] = id; + data_cost_selected[i * disp_step] = minimum; + selected_disparity[i * disp_step] = id; - data_cost[id * cdisp_step1] = numeric_limits::max(); + data_cost[id * disp_step] = numeric_limits::max(); } } } template - __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, int h, int w, int level, int ndisp, float data_weight, float max_data_term, int min_disp) + __global__ void init_data_cost(const uchar *cleft, const uchar *cright, uchar *ctemp, size_t cimg_step, + int h, int w, int level, int ndisp, float data_weight, float max_data_term, + int min_disp, size_t msg_step, size_t disp_step) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -195,7 +199,7 @@ namespace cv { namespace cuda { namespace device int x0 = x << level; int xt = (x + 1) << level; - T* data_cost = (T*)ctemp + y * cmsg_step + x; + T* data_cost = (T*)ctemp + y * msg_step + x; for(int d = 0; d < ndisp; ++d) { @@ -216,13 +220,15 @@ namespace cv { namespace cuda { namespace device } } } - data_cost[cdisp_step1 * d] = saturate_cast(val); + data_cost[disp_step * d] = saturate_cast(val); } } } template - __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 ndisp, float data_weight, float max_data_term, int min_disp) + __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 ndisp, float data_weight, float max_data_term, + int min_disp, size_t msg_step, size_t disp_step) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -261,16 +267,16 @@ namespace cv { namespace cuda { namespace device reduce(smem + winsz * threadIdx.z, val, tid, plus()); - T* data_cost = (T*)ctemp + y_out * cmsg_step + x_out; + T* data_cost = (T*)ctemp + y_out * msg_step + x_out; if (tid == 0) - data_cost[cdisp_step1 * d] = saturate_cast(val); + data_cost[disp_step * d] = saturate_cast(val); } } template - 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, float data_weight, float max_data_term, int min_disp, 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, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -280,15 +286,15 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); break; - case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); break; - case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp); break; + case 1: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); break; + case 3: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); break; + case 4: init_data_cost<<>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); 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, size_t cimg_step, int rows, int cols, int h, int w, int level, int ndisp, int channels, float data_weight, float max_data_term, int min_disp, 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, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -299,9 +305,9 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; - case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; - case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp); break; + case 1: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); break; + case 3: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); break; + case 4: init_data_cost_reduce<<>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } @@ -311,7 +317,7 @@ namespace cv { namespace cuda { namespace device int h, int w, int level, int nr_plane, int ndisp, int channels, float data_weight, float max_data_term, int min_disp, bool use_local_init_data_cost, 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, float data_weight, float max_data_term, int min_disp, 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, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step, cudaStream_t stream); static const InitDataCostCaller init_data_cost_callers[] = { @@ -321,10 +327,8 @@ namespace cv { namespace cuda { namespace device }; size_t disp_step = msg_step * h; - 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, cimg_step, rows, cols, h, w, level, ndisp, channels, data_weight, max_data_term, min_disp, stream); + init_data_cost_callers[level](cleft, cright, ctemp, cimg_step, rows, cols, h, w, level, ndisp, channels, data_weight, max_data_term, min_disp, msg_step, disp_step, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) @@ -337,9 +341,9 @@ namespace cv { namespace cuda { namespace device grid.y = divUp(h, threads.y); if (use_local_init_data_cost == true) - get_first_k_initial_local<<>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp); + get_first_k_initial_local<<>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp, msg_step, disp_step); else - get_first_k_initial_global<<>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp); + get_first_k_initial_global<<>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp, msg_step, disp_step); cudaSafeCall( cudaGetLastError() ); @@ -542,7 +546,7 @@ namespace cv { namespace cuda { namespace device int id = 0; for(int j = 0; j < nr_plane2; j++) { - T cur = data_cost_new[j * cdisp_step1]; + T cur = data_cost_new[j * disp_step1]; if(cur < minimum) { minimum = cur; From 85601e03dd345d88f59ebf5f6b81a42fd336ecaa Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 7 Jul 2014 10:28:47 -0700 Subject: [PATCH 18/19] remove constant memory use in compute_data_cost --- modules/cudastereo/src/cuda/stereocsbp.cu | 43 +++++++++++------------ 1 file changed, 20 insertions(+), 23 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index 974b503733..eb371c1881 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -362,7 +362,7 @@ namespace cv { namespace cuda { namespace device /////////////////////////////////////////////////////////////// template - __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, float data_weight, float max_data_term, int min_disp) + __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, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; @@ -375,8 +375,8 @@ namespace cv { namespace cuda { namespace device int x0 = x << level; int xt = (x + 1) << level; - const T* selected_disparity = selected_disp_pyr + y/2 * cmsg_step + x/2; - T* data_cost = data_cost_ + y * cmsg_step + x; + const T* selected_disparity = selected_disp_pyr + y/2 * msg_step + x/2; + T* data_cost = data_cost_ + y * msg_step + x; for(int d = 0; d < nr_plane; d++) { @@ -385,7 +385,7 @@ namespace cv { namespace cuda { namespace device { for(int xi = x0; xi < xt; xi++) { - int sel_disp = selected_disparity[d * cdisp_step2]; + int sel_disp = selected_disparity[d * disp_step2]; int xr = xi - sel_disp; if (xr < 0 || sel_disp < min_disp) @@ -399,13 +399,13 @@ namespace cv { namespace cuda { namespace device } } } - data_cost[cdisp_step1 * d] = saturate_cast(val); + data_cost[disp_step1 * d] = saturate_cast(val); } } } template - __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, float data_weight, float max_data_term, int min_disp) + __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, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2) { int x_out = blockIdx.x; int y_out = blockIdx.y % h; @@ -413,12 +413,12 @@ namespace cv { namespace cuda { namespace device int tid = threadIdx.x; - const T* selected_disparity = selected_disp_pyr + y_out/2 * cmsg_step + x_out/2; - T* data_cost = data_cost_ + y_out * cmsg_step + x_out; + const T* selected_disparity = selected_disp_pyr + y_out/2 * msg_step + x_out/2; + T* data_cost = data_cost_ + y_out * msg_step + x_out; if (d < nr_plane) { - int sel_disp = selected_disparity[d * cdisp_step2]; + int sel_disp = selected_disparity[d * disp_step2]; int x0 = x_out << level; int y0 = y_out << level; @@ -450,13 +450,13 @@ namespace cv { namespace cuda { namespace device reduce(smem + winsz * threadIdx.z, val, tid, plus()); if (tid == 0) - data_cost[cdisp_step1 * d] = saturate_cast(val); + data_cost[disp_step1 * d] = saturate_cast(val); } } template 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -466,16 +466,16 @@ namespace cv { namespace cuda { namespace device switch(channels) { - case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; - case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; - case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp); break; + case 1: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; + case 3: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; + case 4: compute_data_cost<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } template 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream) + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2, cudaStream_t stream) { const int threadsNum = 256; const size_t smem_size = threadsNum * sizeof(float); @@ -486,9 +486,9 @@ namespace cv { namespace cuda { namespace device switch (channels) { - case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; - case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; - case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp); break; + case 1: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; + case 3: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; + case 4: compute_data_cost_reduce<<>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2); break; default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); } } @@ -499,7 +499,7 @@ namespace cv { namespace cuda { namespace device int min_disp, cudaStream_t stream) { 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, float data_weight, float max_data_term, int min_disp, cudaStream_t stream); + int h, int w, int level, int nr_plane, int channels, float data_weight, float max_data_term, int min_disp, size_t msg_step, size_t disp_step1, size_t disp_step2, cudaStream_t stream); static const ComputeDataCostCaller callers[] = { @@ -510,11 +510,8 @@ namespace cv { namespace cuda { namespace device size_t disp_step1 = msg_step * h; size_t disp_step2 = msg_step * h2; - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step1, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cdisp_step2, &disp_step2, sizeof(size_t)) ); - cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, sizeof(size_t)) ); - callers[level](cleft, cright, cimg_step, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, data_weight, max_data_term, min_disp, stream); + callers[level](cleft, cright, cimg_step, disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, data_weight, max_data_term, min_disp, msg_step, disp_step1, disp_step2, stream); cudaSafeCall( cudaGetLastError() ); if (stream == 0) From 4644689d5a31f0d80527e511b8f86fa05bc8681f Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Thu, 10 Jul 2014 23:53:44 -0700 Subject: [PATCH 19/19] And remove final vestiges. --- modules/cudastereo/src/cuda/stereocsbp.cu | 8 -------- 1 file changed, 8 deletions(-) diff --git a/modules/cudastereo/src/cuda/stereocsbp.cu b/modules/cudastereo/src/cuda/stereocsbp.cu index eb371c1881..dd535e8b20 100644 --- a/modules/cudastereo/src/cuda/stereocsbp.cu +++ b/modules/cudastereo/src/cuda/stereocsbp.cu @@ -54,14 +54,6 @@ namespace cv { namespace cuda { namespace device { namespace stereocsbp { - /////////////////////////////////////////////////////////////// - /////////////////////// load constants //////////////////////// - /////////////////////////////////////////////////////////////// - - __constant__ size_t cmsg_step; - __constant__ size_t cdisp_step1; - __constant__ size_t cdisp_step2; - /////////////////////////////////////////////////////////////// /////////////////////// init data cost //////////////////////// ///////////////////////////////////////////////////////////////