Skip to content

O
opencv
Commit 2832cfdf authored by Aaron Denney's avatar Aaron Denney
Browse files
Options

No longer use constant memory for image step.

parent d8d946a4
Hide whitespace changes
Inline Side-by-side
Showing with 39 additions and 44 deletions
modules/cudastereo/src/cuda/stereocsbp.cu
View file @ 2832cfdf
...@@ -67,14 +67,12 @@ namespace cv { namespace cuda { namespace device ...@@ -67,14 +67,12 @@ namespace cv { namespace cuda { namespace device
__constant__ int cth; __constant__ int cth;
__constant__ size_t cimg_step;
__constant__ size_t cmsg_step; __constant__ size_t cmsg_step;
__constant__ size_t cdisp_step1; __constant__ size_t cdisp_step1;
__constant__ size_t cdisp_step2; __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(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th)
size_t leftstep)
{ {
cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) ); cudaSafeCall( cudaMemcpyToSymbol(cndisp, &ndisp, sizeof(int)) );
...@@ -84,8 +82,6 @@ namespace cv { namespace cuda { namespace device ...@@ -84,8 +82,6 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) );
cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) ); 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 ...@@ -217,7 +213,7 @@ namespace cv { namespace cuda { namespace device
} }
template <typename T, int channels> template <typename T, int channels>
__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 x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y; int y = blockIdx.y * blockDim.y + threadIdx.y;
...@@ -257,7 +253,7 @@ namespace cv { namespace cuda { namespace device ...@@ -257,7 +253,7 @@ namespace cv { namespace cuda { namespace device
} }
template <typename T, int winsz, int channels> template <typename T, int winsz, int channels>
__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 x_out = blockIdx.x;
int y_out = blockIdx.y % h; int y_out = blockIdx.y % h;
...@@ -305,7 +301,7 @@ namespace cv { namespace cuda { namespace device ...@@ -305,7 +301,7 @@ namespace cv { namespace cuda { namespace device
template <typename T> template <typename T>
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 threads(32, 8, 1);
dim3 grid(1, 1, 1); dim3 grid(1, 1, 1);
...@@ -315,15 +311,15 @@ namespace cv { namespace cuda { namespace device ...@@ -315,15 +311,15 @@ namespace cv { namespace cuda { namespace device
switch (channels) switch (channels)
{ {
case 1: init_data_cost<T, 1><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, h, w, level); break; case 1: init_data_cost<T, 1><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level); break;
case 3: init_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, h, w, level); break; case 3: init_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level); break;
case 4: init_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, h, w, level); break; case 4: init_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
} }
} }
template <typename T, int winsz> template <typename T, int winsz>
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 int threadsNum = 256;
const size_t smem_size = threadsNum * sizeof(float); const size_t smem_size = threadsNum * sizeof(float);
...@@ -334,19 +330,19 @@ namespace cv { namespace cuda { namespace device ...@@ -334,19 +330,19 @@ namespace cv { namespace cuda { namespace device
switch (channels) switch (channels)
{ {
case 1: init_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, level, rows, cols, h); break; case 1: init_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break;
case 3: init_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, level, rows, cols, h); break; case 3: init_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break;
case 4: init_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, level, rows, cols, h); break; case 4: init_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
} }
} }
template<class T> template<class T>
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) 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[] = static const InitDataCostCaller init_data_cost_callers[] =
{ {
...@@ -359,7 +355,7 @@ namespace cv { namespace cuda { namespace device ...@@ -359,7 +355,7 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) ); cudaSafeCall( cudaMemcpyToSymbol(cdisp_step1, &disp_step, sizeof(size_t)) );
cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_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() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)
...@@ -382,10 +378,10 @@ namespace cv { namespace cuda { namespace device ...@@ -382,10 +378,10 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaDeviceSynchronize() ); 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); 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); 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 ...@@ -393,7 +389,7 @@ namespace cv { namespace cuda { namespace device
/////////////////////////////////////////////////////////////// ///////////////////////////////////////////////////////////////
template <typename T, int channels> template <typename T, int channels>
__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 x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y; int y = blockIdx.y * blockDim.y + threadIdx.y;
...@@ -436,7 +432,7 @@ namespace cv { namespace cuda { namespace device ...@@ -436,7 +432,7 @@ namespace cv { namespace cuda { namespace device
} }
template <typename T, int winsz, int channels> template <typename T, int winsz, int channels>
__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 x_out = blockIdx.x;
int y_out = blockIdx.y % h; int y_out = blockIdx.y % h;
...@@ -486,7 +482,7 @@ namespace cv { namespace cuda { namespace device ...@@ -486,7 +482,7 @@ namespace cv { namespace cuda { namespace device
} }
template <typename T> template <typename T>
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) int h, int w, int level, int nr_plane, int channels, cudaStream_t stream)
{ {
dim3 threads(32, 8, 1); dim3 threads(32, 8, 1);
...@@ -497,15 +493,15 @@ namespace cv { namespace cuda { namespace device ...@@ -497,15 +493,15 @@ namespace cv { namespace cuda { namespace device
switch(channels) switch(channels)
{ {
case 1: compute_data_cost<T, 1><<<grid, threads, 0, stream>>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; case 1: compute_data_cost<T, 1><<<grid, threads, 0, stream>>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break;
case 3: compute_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; case 3: compute_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break;
case 4: compute_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, disp_selected_pyr, data_cost, h, w, level, nr_plane); break; case 4: compute_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count"); default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
} }
} }
template <typename T, int winsz> template <typename T, int winsz>
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) int h, int w, int level, int nr_plane, int channels, cudaStream_t stream)
{ {
const int threadsNum = 256; const int threadsNum = 256;
...@@ -517,18 +513,18 @@ namespace cv { namespace cuda { namespace device ...@@ -517,18 +513,18 @@ namespace cv { namespace cuda { namespace device
switch (channels) switch (channels)
{ {
case 1: compute_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; case 1: compute_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break;
case 3: compute_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; case 3: compute_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, cimg_step, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break;
case 4: compute_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(cleft, cright, disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break; case 4: compute_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(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"); default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
} }
} }
template<class T> template<class T>
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) 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); int h, int w, int level, int nr_plane, int channels, cudaStream_t stream);
static const ComputeDataCostCaller callers[] = static const ComputeDataCostCaller callers[] =
...@@ -544,17 +540,17 @@ namespace cv { namespace cuda { namespace device ...@@ -544,17 +540,17 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaMemcpyToSymbol(cdisp_step2, &disp_step2, sizeof(size_t)) ); cudaSafeCall( cudaMemcpyToSymbol(cdisp_step2, &disp_step2, sizeof(size_t)) );
cudaSafeCall( cudaMemcpyToSymbol(cmsg_step, &msg_step, 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() ); cudaSafeCall( cudaGetLastError() );
if (stream == 0) if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() ); 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); 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); int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream);
......
modules/cudastereo/src/cuda/stereocsbp.hpp
View file @ 2832cfdf
...@@ -2,15 +2,14 @@ namespace cv { namespace cuda { namespace device ...@@ -2,15 +2,14 @@ namespace cv { namespace cuda { namespace device
{ {
namespace stereocsbp 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(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th);
size_t leftstep);
template<class T> template<class T>
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); int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream);
template<class T> template<class T>
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); int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream);
template<class T> template<class T>
......
modules/cudastereo/src/stereocsbp.cpp
View file @ 2832cfdf
...@@ -222,7 +222,7 @@ namespace ...@@ -222,7 +222,7 @@ namespace
//////////////////////////////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////
// Compute // 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); l[0].setTo(0, _stream);
d[0].setTo(0, _stream); d[0].setTo(0, _stream);
...@@ -245,12 +245,12 @@ namespace ...@@ -245,12 +245,12 @@ namespace
{ {
if (i == levels_ - 1) if (i == levels_ - 1)
{ {
init_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), temp_.ptr<uchar>(), left.rows, left.cols, disp_selected_pyr[cur_idx].ptr<float>(), data_cost_selected.ptr<float>(), init_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), temp_.ptr<uchar>(), left.step, left.rows, left.cols, disp_selected_pyr[cur_idx].ptr<float>(), data_cost_selected.ptr<float>(),
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(), use_local_init_data_cost_, stream);
} }
else else
{ {
compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), disp_selected_pyr[cur_idx].ptr<float>(), data_cost.ptr<float>(), elem_step, compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), left.step, disp_selected_pyr[cur_idx].ptr<float>(), data_cost.ptr<float>(), 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(), stream);
int new_idx = (cur_idx + 1) & 1; int new_idx = (cur_idx + 1) & 1;
...@@ -276,12 +276,12 @@ namespace ...@@ -276,12 +276,12 @@ namespace
{ {
if (i == levels_ - 1) if (i == levels_ - 1)
{ {
init_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), temp_.ptr<uchar>(), left.rows, left.cols, disp_selected_pyr[cur_idx].ptr<short>(), data_cost_selected.ptr<short>(), init_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), temp_.ptr<uchar>(), left.step, left.rows, left.cols, disp_selected_pyr[cur_idx].ptr<short>(), data_cost_selected.ptr<short>(),
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(), use_local_init_data_cost_, stream);
} }
else else
{ {
compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), disp_selected_pyr[cur_idx].ptr<short>(), data_cost.ptr<short>(), elem_step, compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), left.step, disp_selected_pyr[cur_idx].ptr<short>(), data_cost.ptr<short>(), 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(), stream);
int new_idx = (cur_idx + 1) & 1; int new_idx = (cur_idx + 1) & 1;
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment