generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Merge pull request #2983 from wnoise:shrink-global-cuda-usage

Browse Source
This commit is contained in:
Vadim Pisarevsky 2014-07-30 15:55:02 +00:00
commit 1a1097ab23
6 changed files with 267 additions and 337 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

46
modules/cudastereo/src/cuda/disparity_bilateral_filter.cu
View File

@ -45,34 +45,12 @@
#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
{
__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 <int channels>
struct DistRgbMax
{
@ -95,7 +73,11 @@ namespace cv { namespace cuda { namespace device
};
template <int channels, typename T>
__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 +160,7 @@ namespace cv { namespace cuda { namespace device
}
template <typename T>
void disp_bilateral_filter(PtrStepSz<T> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream)
void disp_bilateral_filter(PtrStepSz<T> 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 +172,20 @@ namespace cv { namespace cuda { namespace device
case 1:
for (int i = 0; i < iters; ++i)
{
disp_bilateral_filter<1><<<grid, threads, 0, stream>>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols);
disp_bilateral_filter<1><<<grid, threads, 0, stream>>>(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><<<grid, threads, 0, stream>>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols);
disp_bilateral_filter<1><<<grid, threads, 0, stream>>>(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><<<grid, threads, 0, stream>>>(0, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols);
disp_bilateral_filter<3><<<grid, threads, 0, stream>>>(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><<<grid, threads, 0, stream>>>(1, disp.data, disp.step/sizeof(T), img.data, img.step, disp.rows, disp.cols);
disp_bilateral_filter<3><<<grid, threads, 0, stream>>>(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 +197,8 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaDeviceSynchronize() );
}
template void disp_bilateral_filter<uchar>(PtrStepSz<uchar> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream);
template void disp_bilateral_filter<short>(PtrStepSz<short> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream);
template void disp_bilateral_filter<uchar>(PtrStepSz<uchar> 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<short>(PtrStepSz<short> 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

8
modules/cudastereo/src/cuda/disparity_bilateral_filter.hpp Normal file
View File

@ -0,0 +1,8 @@
namespace cv { namespace cuda { namespace device
{
namespace disp_bilateral_filter
{
template<typename T>
void disp_bilateral_filter(PtrStepSz<T> disp, PtrStepSzb img, int channels, int iters, const float *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream);
}
}}}

442
modules/cudastereo/src/cuda/stereocsbp.cu
View File

@ -48,109 +48,61 @@
#include "opencv2/core/cuda/reduce.hpp"
#include "opencv2/core/cuda/functional.hpp"
#include "cuda/stereocsbp.hpp"
namespace cv { namespace cuda { namespace device
{
namespace stereocsbp
{
///////////////////////////////////////////////////////////////
/////////////////////// load constants ////////////////////////
///////////////////////////////////////////////////////////////
__constant__ int cndisp;
__constant__ float cmax_data_term;
__constant__ float cdata_weight;
__constant__ float cmax_disc_term;
__constant__ float cdisc_single_jump;
__constant__ int cth;
__constant__ size_t cimg_step;
__constant__ size_t cmsg_step;
__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)
{
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)) );
cudaSafeCall( cudaMemcpyToSymbol(cdisc_single_jump, &disc_single_jump, sizeof(float)) );
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)) );
}
///////////////////////////////////////////////////////////////
/////////////////////// init data cost ////////////////////////
///////////////////////////////////////////////////////////////
template <int channels> struct DataCostPerPixel;
template <> struct DataCostPerPixel<1>
template <int channels> 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)
{
static __device__ __forceinline__ float compute(const uchar* left, const uchar* right)
{
return fmin(cdata_weight * ::abs((int)*left - *right), cdata_weight * cmax_data_term);
}
};
template <> struct DataCostPerPixel<3>
return fmin( ::abs((int)*left - *right), max_data_term);
}
template<> __device__ __forceinline__ static float pixeldiff<3>(const uchar* left, const uchar* right, float max_data_term)
{
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 fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term);
}
};
template <> struct DataCostPerPixel<4>
return fmin(tr + tg + tb, max_data_term);
}
template<> __device__ __forceinline__ static float pixeldiff<4>(const uchar* left, const uchar* right, float max_data_term)
{
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 fmin(cdata_weight * (tr + tg + tb), cdata_weight * cmax_data_term);
}
};
return fmin(tr + tg + tb, max_data_term);
}
template <typename T>
__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 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++)
{
T minimum = device::numeric_limits<T>::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];
T cur = data_cost[d * disp_step];
if(cur < minimum)
{
minimum = cur;
@ -158,46 +110,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<T>::max();
data_cost_selected[i * disp_step] = minimum;
selected_disparity[i * disp_step] = id;
data_cost [id * disp_step] = numeric_limits<T>::max();
}
}
}
template <typename T>
__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 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 < 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)
{
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<T>::max();
data_cost[d * disp_step] = numeric_limits<T>::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++)
@ -205,25 +158,27 @@ namespace cv { namespace cuda { namespace device
T minimum = numeric_limits<T>::max();
int id = 0;
for (int d = 0; d < cndisp; d++)
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<T>::max();
data_cost[id * disp_step] = numeric_limits<T>::max();
}
}
}
template <typename T, int channels>
__global__ void init_data_cost(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, 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;
@ -236,9 +191,9 @@ 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 < cndisp; ++d)
for(int d = 0; d < ndisp; ++d)
{
float val = 0.0f;
for(int yi = y0; yi < yt; yi++)
@ -246,24 +201,26 @@ namespace cv { namespace cuda { namespace device
for(int xi = x0; xi < xt; xi++)
{
int xr = xi - d;
if(d < cth || xr < 0)
val += cdata_weight * cmax_data_term;
if(d < min_disp || xr < 0)
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 += DataCostPerPixel<channels>::compute(lle, lri);
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
}
}
}
data_cost[cdisp_step1 * d] = saturate_cast<T>(val);
data_cost[disp_step * d] = saturate_cast<T>(val);
}
}
}
template <typename T, int winsz, int channels>
__global__ void init_data_cost_reduce(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, 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;
@ -271,7 +228,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;
@ -281,8 +238,8 @@ namespace cv { namespace cuda { namespace device
float val = 0.0f;
if (x0 + tid < cols)
{
if (x0 + tid - d < 0 || d < cth)
val = cdata_weight * cmax_data_term * len;
if (x0 + tid - d < 0 || d < min_disp)
val = data_weight * max_data_term * len;
else
{
const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid );
@ -290,7 +247,7 @@ namespace cv { namespace cuda { namespace device
for(int y = 0; y < len; ++y)
{
val += DataCostPerPixel<channels>::compute(lle, lri);
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
lle += cimg_step;
lri += cimg_step;
@ -302,16 +259,16 @@ namespace cv { namespace cuda { namespace device
reduce<winsz>(smem + winsz * threadIdx.z, val, tid, plus<float>());
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<T>(val);
data_cost[disp_step * d] = saturate_cast<T>(val);
}
}
template <typename T>
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_(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);
@ -321,15 +278,15 @@ namespace cv { namespace cuda { namespace device
switch (channels)
{
case 1: init_data_cost<T, 1><<<grid, threads, 0, stream>>>(h, w, level); break;
case 3: init_data_cost<T, 3><<<grid, threads, 0, stream>>>(h, w, level); break;
case 4: init_data_cost<T, 4><<<grid, threads, 0, stream>>>(h, w, level); 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, min_disp, msg_step, disp_step); 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, min_disp, msg_step, disp_step); 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, min_disp, msg_step, disp_step); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
}
}
template <typename T, int winsz>
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_(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);
@ -340,19 +297,19 @@ namespace cv { namespace cuda { namespace device
switch (channels)
{
case 1: init_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(level, rows, cols, h); break;
case 3: init_data_cost_reduce<T, winsz, 3><<<grid, threads, smem_size, stream>>>(level, rows, cols, h); break;
case 4: init_data_cost_reduce<T, winsz, 4><<<grid, threads, smem_size, stream>>>(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, ndisp, data_weight, max_data_term, min_disp, msg_step, disp_step); 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, min_disp, msg_step, disp_step); 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, min_disp, msg_step, disp_step); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
}
}
template<class T>
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)
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, int min_disp, 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)(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[] =
{
@ -362,10 +319,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](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, min_disp, msg_step, disp_step, stream);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)
@ -378,9 +333,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<<<grid, threads, 0, stream>>> (data_cost_selected, disp_selected_pyr, h, w, nr_plane);
get_first_k_initial_local<<<grid, threads, 0, stream>>> (ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp, msg_step, disp_step);
else
get_first_k_initial_global<<<grid, threads, 0, stream>>>(data_cost_selected, disp_selected_pyr, h, w, nr_plane);
get_first_k_initial_global<<<grid, threads, 0, stream>>>(ctemp, data_cost_selected, disp_selected_pyr, h, w, nr_plane, ndisp, msg_step, disp_step);
cudaSafeCall( cudaGetLastError() );
@ -388,18 +343,18 @@ 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,
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<short>(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(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);
template void init_data_cost<float>(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 <typename T, int channels>
__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(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;
@ -412,8 +367,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++)
{
@ -422,27 +377,27 @@ 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 < cth)
val += cdata_weight * cmax_data_term;
if (xr < 0 || sel_disp < min_disp)
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 += DataCostPerPixel<channels>::compute(left_x, right_x);
val += data_weight * pixeldiff<channels>(left_x, right_x, max_data_term);
}
}
}
data_cost[cdisp_step1 * d] = saturate_cast<T>(val);
data_cost[disp_step1 * d] = saturate_cast<T>(val);
}
}
}
template <typename T, int winsz, int channels>
__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(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;
@ -450,12 +405,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;
@ -465,8 +420,8 @@ namespace cv { namespace cuda { namespace device
float val = 0.0f;
if (x0 + tid < cols)
{
if (x0 + tid - sel_disp < 0 || sel_disp < cth)
val = cdata_weight * cmax_data_term * len;
if (x0 + tid - sel_disp < 0 || sel_disp < min_disp)
val = data_weight * max_data_term * len;
else
{
const uchar* lle = cleft + y0 * cimg_step + channels * (x0 + tid );
@ -474,7 +429,7 @@ namespace cv { namespace cuda { namespace device
for(int y = 0; y < len; ++y)
{
val += DataCostPerPixel<channels>::compute(lle, lri);
val += data_weight * pixeldiff<channels>(lle, lri, max_data_term);
lle += cimg_step;
lri += cimg_step;
@ -487,13 +442,13 @@ namespace cv { namespace cuda { namespace device
reduce<winsz>(smem + winsz * threadIdx.z, val, tid, plus<float>());
if (tid == 0)
data_cost[cdisp_step1 * d] = saturate_cast<T>(val);
data_cost[disp_step1 * d] = saturate_cast<T>(val);
}
}
template <typename T>
void compute_data_cost_caller_(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)
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, 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);
@ -503,16 +458,16 @@ namespace cv { namespace cuda { namespace device
switch(channels)
{
case 1: compute_data_cost<T, 1><<<grid, threads, 0, stream>>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break;
case 3: compute_data_cost<T, 3><<<grid, threads, 0, stream>>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break;
case 4: compute_data_cost<T, 4><<<grid, threads, 0, stream>>>(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, min_disp, msg_step, disp_step1, disp_step2); 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, min_disp, msg_step, disp_step1, disp_step2); 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, min_disp, msg_step, disp_step1, disp_step2); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
}
}
template <typename T, int winsz>
void compute_data_cost_reduce_caller_(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)
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, 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);
@ -523,19 +478,20 @@ namespace cv { namespace cuda { namespace device
switch (channels)
{
case 1: compute_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(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>>>(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>>>(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, min_disp, msg_step, disp_step1, disp_step2); 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, min_disp, msg_step, disp_step1, disp_step2); 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, min_disp, msg_step, disp_step1, disp_step2); break;
default: CV_Error(cv::Error::BadNumChannels, "Unsupported channels count");
}
}
template<class T>
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)
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,
int min_disp, cudaStream_t stream)
{
typedef void (*ComputeDataCostCaller)(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);
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, size_t msg_step, size_t disp_step1, size_t disp_step2, cudaStream_t stream);
static const ComputeDataCostCaller callers[] =
{
@ -546,22 +502,19 @@ 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](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, min_disp, msg_step, disp_step1, disp_step2, 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,
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, 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, int min_disp, cudaStream_t stream);
template void compute_data_cost(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);
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, int min_disp, cudaStream_t stream);
///////////////////////////////////////////////////////////////
@ -574,7 +527,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++)
{
@ -582,7 +535,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;
@ -590,70 +543,72 @@ 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<T>::max();
data_cost_new[id * disp_step1] = numeric_limits<T>::max();
}
}
template <typename T>
__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_,
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);
}
}
template<class T>
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,
@ -662,9 +617,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);
@ -672,11 +624,12 @@ namespace cv { namespace cuda { namespace device
grid.x = divUp(w, threads.x);
grid.y = divUp(h, threads.y);
init_message<<<grid, threads, 0, stream>>>(u_new, d_new, l_new, r_new,
init_message<<<grid, threads, 0, stream>>>(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,
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)
@ -684,13 +637,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,
@ -702,13 +655,14 @@ namespace cv { namespace cuda { namespace device
template <typename T>
__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, float disc_single_jump, volatile T* temp,
size_t disp_step)
{
T minimum = numeric_limits<T>::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)
@ -720,55 +674,53 @@ namespace cv { namespace cuda { namespace device
float sum = 0;
for(int d = 0; d < nr_plane; d++)
{
float cost_min = minimum + cmax_disc_term;
T src_disp_reg = src_disp[d * cdisp_step1];
float cost_min = minimum + max_disc_term;
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] + cdisc_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<T>(cost_min);
temp[d * disp_step] = saturate_cast<T>(cost_min);
sum += cost_min;
}
sum /= nr_plane;
for(int d = 0; d < nr_plane; d++)
msg_dst[d * cdisp_step1] = saturate_cast<T>(temp[d * cdisp_step1] - sum);
msg_dst[d * disp_step] = saturate_cast<T>(temp[d * disp_step] - sum);
}
template <typename T>
__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 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, 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 + 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);
}
}
template<class T>
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)
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, 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);
@ -778,18 +730,18 @@ namespace cv { namespace cuda { namespace device
for(int t = 0; t < iters; ++t)
{
compute_message<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1);
compute_message<<<grid, threads, 0, stream>>>(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)
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,
int h, int w, int nr_plane, int iters, cudaStream_t stream);
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, float disc_single_jump, 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,
int h, int w, int nr_plane, int iters, 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, float disc_single_jump, cudaStream_t stream);
///////////////////////////////////////////////////////////////
@ -800,26 +752,26 @@ namespace cv { namespace cuda { namespace device
template <typename T>
__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<short> disp, int nr_plane)
PtrStepSz<short> 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<T>::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)
@ -837,8 +789,6 @@ namespace cv { namespace cuda { namespace device
const PtrStepSz<short>& 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);
@ -846,7 +796,7 @@ namespace cv { namespace cuda { namespace device
grid.x = divUp(disp.cols, threads.x);
grid.y = divUp(disp.rows, threads.y);
compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected, disp, nr_plane);
compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected, disp, nr_plane, msg_step, disp_step);
cudaSafeCall( cudaGetLastError() );
if (stream == 0)

29
modules/cudastereo/src/cuda/stereocsbp.hpp Normal file
View File

@ -0,0 +1,29 @@
namespace cv { namespace cuda { namespace device
{
namespace stereocsbp
{
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,
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<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,
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<class T>
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<class T>
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, float disc_single_jump, cudaStream_t stream);
template<class T>
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<short>& disp, int nr_plane, cudaStream_t stream);
}
}}}

15
modules/cudastereo/src/disparity_bilateral_filter.cpp
View File

@ -51,16 +51,7 @@ Ptr<cuda::DisparityBilateralFilter> cv::cuda::createDisparityBilateralFilter(int
#else /* !defined (HAVE_CUDA) */
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<typename T>
void disp_bilateral_filter(PtrStepSz<T> disp, PtrStepSzb img, int channels, int iters, cudaStream_t stream);
}
}}}
#include "cuda/disparity_bilateral_filter.hpp"
namespace
{
@ -165,7 +156,7 @@ namespace
const short edge_disc = std::max<short>(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<float>(), 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 +164,7 @@ namespace
if (dst.data != disp.data)
disp.copyTo(dst, stream);
disp_bilateral_filter<T>(dst, img, img.channels(), iters, StreamAccessor::getStream(stream));
disp_bilateral_filter<T>(dst, img, img.channels(), iters, table_color.ptr<float>(), (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)

64
modules/cudastereo/src/stereocsbp.cpp
View File

@ -53,37 +53,7 @@ Ptr<cuda::StereoConstantSpaceBP> 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<class T>
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<class T>
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<class T>
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<class T>
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<class T>
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<short>& disp, int nr_plane, cudaStream_t stream);
}
}}}
#include "cuda/stereocsbp.hpp"
namespace
{
@ -252,8 +222,6 @@ namespace
////////////////////////////////////////////////////////////////////////////
// Compute
load_constants(ndisp_, max_data_term_, data_weight_, max_disc_term_, disc_single_jump_, min_disp_th_, left, right, temp_);
l[0].setTo(0, _stream);
d[0].setTo(0, _stream);
r[0].setTo(0, _stream);
@ -275,17 +243,18 @@ namespace
{
if (i == levels_ - 1)
{
init_data_cost(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);
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(), data_weight_, max_data_term_, min_disp_th_, use_local_init_data_cost_, stream);
}
else
{
compute_data_cost(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);
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(), data_weight_, max_data_term_, min_disp_th_, stream);
int new_idx = (cur_idx + 1) & 1;
init_message(u[new_idx].ptr<float>(), d[new_idx].ptr<float>(), l[new_idx].ptr<float>(), r[new_idx].ptr<float>(),
init_message(temp_.ptr<uchar>(),
u[new_idx].ptr<float>(), d[new_idx].ptr<float>(), l[new_idx].ptr<float>(), r[new_idx].ptr<float>(),
u[cur_idx].ptr<float>(), d[cur_idx].ptr<float>(), l[cur_idx].ptr<float>(), r[cur_idx].ptr<float>(),
disp_selected_pyr[new_idx].ptr<float>(), disp_selected_pyr[cur_idx].ptr<float>(),
data_cost_selected.ptr<float>(), data_cost.ptr<float>(), elem_step, rows_pyr[i],
@ -294,9 +263,9 @@ namespace
cur_idx = new_idx;
}
calc_all_iterations(u[cur_idx].ptr<float>(), d[cur_idx].ptr<float>(), l[cur_idx].ptr<float>(), r[cur_idx].ptr<float>(),
calc_all_iterations(temp_.ptr<uchar>(), u[cur_idx].ptr<float>(), d[cur_idx].ptr<float>(), l[cur_idx].ptr<float>(), r[cur_idx].ptr<float>(),
data_cost_selected.ptr<float>(), disp_selected_pyr[cur_idx].ptr<float>(), 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_, disc_single_jump_, stream);
}
}
else
@ -305,17 +274,18 @@ namespace
{
if (i == levels_ - 1)
{
init_data_cost(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);
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(), data_weight_, max_data_term_, min_disp_th_, use_local_init_data_cost_, stream);
}
else
{
compute_data_cost(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);
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(), data_weight_, max_data_term_, min_disp_th_, stream);
int new_idx = (cur_idx + 1) & 1;
init_message(u[new_idx].ptr<short>(), d[new_idx].ptr<short>(), l[new_idx].ptr<short>(), r[new_idx].ptr<short>(),
init_message(temp_.ptr<uchar>(),
u[new_idx].ptr<short>(), d[new_idx].ptr<short>(), l[new_idx].ptr<short>(), r[new_idx].ptr<short>(),
u[cur_idx].ptr<short>(), d[cur_idx].ptr<short>(), l[cur_idx].ptr<short>(), r[cur_idx].ptr<short>(),
disp_selected_pyr[new_idx].ptr<short>(), disp_selected_pyr[cur_idx].ptr<short>(),
data_cost_selected.ptr<short>(), data_cost.ptr<short>(), elem_step, rows_pyr[i],
@ -324,9 +294,9 @@ namespace
cur_idx = new_idx;
}
calc_all_iterations(u[cur_idx].ptr<short>(), d[cur_idx].ptr<short>(), l[cur_idx].ptr<short>(), r[cur_idx].ptr<short>(),
calc_all_iterations(temp_.ptr<uchar>(), u[cur_idx].ptr<short>(), d[cur_idx].ptr<short>(), l[cur_idx].ptr<short>(), r[cur_idx].ptr<short>(),
data_cost_selected.ptr<short>(), disp_selected_pyr[cur_idx].ptr<short>(), 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_, disc_single_jump_, stream);
}
}