From eed5cbc5db780e2c8ee5293f662cb663742c4e3e Mon Sep 17 00:00:00 2001 From: Aaron Denney Date: Mon, 30 Jun 2014 13:22:13 -0700 Subject: [PATCH] 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 3ef665b..792df0c 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 8022475..04337bb 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 12d6731..14d6d25 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; -- 2.7.4