mirror of
https://github.com/opencv/opencv.git
synced 2025-01-19 15:04:01 +08:00
More constant removal.
This commit is contained in:
parent
3ab117df04
commit
eed5cbc5db
@ -58,8 +58,6 @@ namespace cv { namespace cuda { namespace device
|
|||||||
/////////////////////// load constants ////////////////////////
|
/////////////////////// load constants ////////////////////////
|
||||||
///////////////////////////////////////////////////////////////
|
///////////////////////////////////////////////////////////////
|
||||||
|
|
||||||
__constant__ float cmax_data_term;
|
|
||||||
__constant__ float cdata_weight;
|
|
||||||
__constant__ float cdisc_single_jump;
|
__constant__ float cdisc_single_jump;
|
||||||
|
|
||||||
__constant__ int cth;
|
__constant__ int cth;
|
||||||
@ -69,10 +67,8 @@ namespace cv { namespace cuda { namespace device
|
|||||||
__constant__ size_t cdisp_step2;
|
__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(cdisc_single_jump, &disc_single_jump, sizeof(float)) );
|
||||||
|
|
||||||
cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) );
|
cudaSafeCall( cudaMemcpyToSymbol(cth, &min_disp_th, sizeof(int)) );
|
||||||
@ -82,20 +78,20 @@ namespace cv { namespace cuda { namespace device
|
|||||||
/////////////////////// init data cost ////////////////////////
|
/////////////////////// init data cost ////////////////////////
|
||||||
///////////////////////////////////////////////////////////////
|
///////////////////////////////////////////////////////////////
|
||||||
|
|
||||||
template <int channels> float __device__ pixeldiff(const uchar* left, const uchar* right);
|
template <int channels> static float __device__ pixeldiff(const uchar* left, const uchar* right, float max_data_term);
|
||||||
template<> float __device__ __forceinline__ pixeldiff<1>(const uchar* left, const uchar* right)
|
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 tb = 0.114f * ::abs((int)left[0] - right[0]);
|
||||||
float tg = 0.587f * ::abs((int)left[1] - right[1]);
|
float tg = 0.587f * ::abs((int)left[1] - right[1]);
|
||||||
float tr = 0.299f * ::abs((int)left[2] - right[2]);
|
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 l = *((const uchar4*)left);
|
||||||
uchar4 r = *((const uchar4*)right);
|
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 tg = 0.587f * ::abs((int)l.y - r.y);
|
||||||
float tr = 0.299f * ::abs((int)l.z - r.z);
|
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 <typename T>
|
template <typename T>
|
||||||
@ -198,7 +194,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, 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 x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
@ -222,13 +218,13 @@ namespace cv { namespace cuda { namespace device
|
|||||||
{
|
{
|
||||||
int xr = xi - d;
|
int xr = xi - d;
|
||||||
if(d < cth || xr < 0)
|
if(d < cth || xr < 0)
|
||||||
val += cdata_weight * cmax_data_term;
|
val += data_weight * max_data_term;
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
const uchar* lle = cleft + yi * cimg_step + xi * channels;
|
const uchar* lle = cleft + yi * cimg_step + xi * channels;
|
||||||
const uchar* lri = cright + yi * cimg_step + xr * channels;
|
const uchar* lri = cright + yi * cimg_step + xr * channels;
|
||||||
|
|
||||||
val += pixeldiff<channels>(lle, lri);
|
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@ -238,7 +234,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, 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 x_out = blockIdx.x;
|
||||||
int y_out = blockIdx.y % h;
|
int y_out = blockIdx.y % h;
|
||||||
@ -257,7 +253,7 @@ namespace cv { namespace cuda { namespace device
|
|||||||
if (x0 + tid < cols)
|
if (x0 + tid < cols)
|
||||||
{
|
{
|
||||||
if (x0 + tid - d < 0 || d < cth)
|
if (x0 + tid - d < 0 || d < cth)
|
||||||
val = cdata_weight * cmax_data_term * len;
|
val = data_weight * max_data_term * len;
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid );
|
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)
|
for(int y = 0; y < len; ++y)
|
||||||
{
|
{
|
||||||
val += pixeldiff<channels>(lle, lri);
|
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
|
||||||
|
|
||||||
lle += cimg_step;
|
lle += cimg_step;
|
||||||
lri += cimg_step;
|
lri += cimg_step;
|
||||||
@ -286,7 +282,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, 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 threads(32, 8, 1);
|
||||||
dim3 grid(1, 1, 1);
|
dim3 grid(1, 1, 1);
|
||||||
@ -296,15 +292,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, cimg_step, h, w, level, ndisp); break;
|
case 1: init_data_cost<T, 1><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break;
|
||||||
case 3: init_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break;
|
case 3: init_data_cost<T, 3><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); break;
|
||||||
case 4: init_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp); break;
|
case 4: init_data_cost<T, 4><<<grid, threads, 0, stream>>>(cleft, cright, ctemp, cimg_step, h, w, level, ndisp, data_weight, max_data_term); 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, 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 int threadsNum = 256;
|
||||||
const size_t smem_size = threadsNum * sizeof(float);
|
const size_t smem_size = threadsNum * sizeof(float);
|
||||||
@ -315,19 +311,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, cimg_step, level, rows, cols, h, ndisp); break;
|
case 1: init_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break;
|
||||||
case 3: init_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break;
|
case 3: init_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp, data_weight, max_data_term); break;
|
||||||
case 4: init_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(cleft, cright, ctemp, cimg_step, level, rows, cols, h, ndisp); break;
|
case 4: init_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(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");
|
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, size_t cimg_step, 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, 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[] =
|
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(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, 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() );
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
if (stream == 0)
|
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,
|
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,
|
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 //////////////////////
|
////////////////////// compute data cost //////////////////////
|
||||||
///////////////////////////////////////////////////////////////
|
///////////////////////////////////////////////////////////////
|
||||||
|
|
||||||
template <typename T, int channels>
|
template <typename T, int channels>
|
||||||
__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 x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||||
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||||
@ -401,13 +397,13 @@ namespace cv { namespace cuda { namespace device
|
|||||||
int xr = xi - sel_disp;
|
int xr = xi - sel_disp;
|
||||||
|
|
||||||
if (xr < 0 || sel_disp < cth)
|
if (xr < 0 || sel_disp < cth)
|
||||||
val += cdata_weight * cmax_data_term;
|
val += data_weight * max_data_term;
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
const uchar* left_x = cleft + yi * cimg_step + xi * channels;
|
const uchar* left_x = cleft + yi * cimg_step + xi * channels;
|
||||||
const uchar* right_x = cright + yi * cimg_step + xr * channels;
|
const uchar* right_x = cright + yi * cimg_step + xr * channels;
|
||||||
|
|
||||||
val += pixeldiff<channels>(left_x, right_x);
|
val += data_weight * pixeldiff<channels>(left_x, right_x, max_data_term);
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
}
|
}
|
||||||
@ -417,7 +413,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, 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 x_out = blockIdx.x;
|
||||||
int y_out = blockIdx.y % h;
|
int y_out = blockIdx.y % h;
|
||||||
@ -441,7 +437,7 @@ namespace cv { namespace cuda { namespace device
|
|||||||
if (x0 + tid < cols)
|
if (x0 + tid < cols)
|
||||||
{
|
{
|
||||||
if (x0 + tid - sel_disp < 0 || sel_disp < cth)
|
if (x0 + tid - sel_disp < 0 || sel_disp < cth)
|
||||||
val = cdata_weight * cmax_data_term * len;
|
val = data_weight * max_data_term * len;
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid );
|
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)
|
for(int y = 0; y < len; ++y)
|
||||||
{
|
{
|
||||||
val += pixeldiff<channels>(lle, lri);
|
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
|
||||||
|
|
||||||
lle += cimg_step;
|
lle += cimg_step;
|
||||||
lri += cimg_step;
|
lri += cimg_step;
|
||||||
@ -468,7 +464,7 @@ namespace cv { namespace cuda { namespace device
|
|||||||
|
|
||||||
template <typename T>
|
template <typename T>
|
||||||
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*/,
|
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 threads(32, 8, 1);
|
||||||
dim3 grid(1, 1, 1);
|
dim3 grid(1, 1, 1);
|
||||||
@ -478,16 +474,16 @@ namespace cv { namespace cuda { namespace device
|
|||||||
|
|
||||||
switch(channels)
|
switch(channels)
|
||||||
{
|
{
|
||||||
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 1: compute_data_cost<T, 1><<<grid, threads, 0, stream>>>(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<T, 3><<<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, cimg_step, disp_selected_pyr, data_cost, h, w, level, nr_plane, data_weight, max_data_term); 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;
|
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, data_weight, max_data_term); 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, size_t cimg_step, 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, float data_weight, float max_data_term, 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);
|
||||||
@ -498,19 +494,19 @@ 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, cimg_step, 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, data_weight, max_data_term); 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 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, data_weight, max_data_term); 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;
|
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, data_weight, max_data_term); 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, size_t cimg_step, 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, 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,
|
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[] =
|
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(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, 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() );
|
cudaSafeCall( cudaGetLastError() );
|
||||||
|
|
||||||
if (stream == 0)
|
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,
|
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,
|
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);
|
||||||
|
|
||||||
|
|
||||||
///////////////////////////////////////////////////////////////
|
///////////////////////////////////////////////////////////////
|
||||||
|
@ -2,15 +2,15 @@ namespace cv { namespace cuda { namespace device
|
|||||||
{
|
{
|
||||||
namespace stereocsbp
|
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<class T>
|
template<class T>
|
||||||
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,
|
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<class T>
|
template<class T>
|
||||||
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,
|
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<class T>
|
template<class T>
|
||||||
void init_message(uchar *ctemp, 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,
|
||||||
|
@ -222,7 +222,7 @@ namespace
|
|||||||
////////////////////////////////////////////////////////////////////////////
|
////////////////////////////////////////////////////////////////////////////
|
||||||
// Compute
|
// 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);
|
l[0].setTo(0, _stream);
|
||||||
d[0].setTo(0, _stream);
|
d[0].setTo(0, _stream);
|
||||||
@ -246,12 +246,12 @@ namespace
|
|||||||
if (i == levels_ - 1)
|
if (i == levels_ - 1)
|
||||||
{
|
{
|
||||||
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>(),
|
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(), data_weight_, max_data_term_, use_local_init_data_cost_, stream);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), left.step, 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(), data_weight_, max_data_term_, stream);
|
||||||
|
|
||||||
int new_idx = (cur_idx + 1) & 1;
|
int new_idx = (cur_idx + 1) & 1;
|
||||||
|
|
||||||
@ -277,12 +277,12 @@ namespace
|
|||||||
if (i == levels_ - 1)
|
if (i == levels_ - 1)
|
||||||
{
|
{
|
||||||
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>(),
|
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(), data_weight_, max_data_term_, use_local_init_data_cost_, stream);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
compute_data_cost(left.ptr<uchar>(), right.ptr<uchar>(), left.step, 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(), data_weight_, max_data_term_, stream);
|
||||||
|
|
||||||
int new_idx = (cur_idx + 1) & 1;
|
int new_idx = (cur_idx + 1) & 1;
|
||||||
|
|
||||||
|
Loading…
Reference in New Issue
Block a user