#define SHRT_MAX 32767\r
#endif\r
\r
-template <typename T> \r
+template <typename T>\r
struct TypeLimits {};\r
\r
-template <> \r
+template <>\r
struct TypeLimits<short>\r
{\r
static __device__ short max() {return SHRT_MAX;}\r
};\r
\r
-template <> \r
+template <>\r
struct TypeLimits<float>\r
{\r
static __device__ float max() {return FLT_MAX;}\r
__constant__ float cdata_weight;\r
__constant__ float cmax_disc_term;\r
__constant__ float cdisc_single_jump;\r
- \r
+\r
__constant__ int cth;\r
\r
__constant__ size_t cimg_step;\r
__constant__ uchar* ctemp;\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th,\r
const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp)\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdata_weight, &data_weight, sizeof(float)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmax_disc_term, &max_disc_term, sizeof(float)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisc_single_jump, &disc_single_jump, sizeof(float)) );\r
- \r
+\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cth, &min_disp_th, sizeof(int)) );\r
- \r
+\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cimg_step, &left.step, sizeof(size_t)) );\r
\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cleft, &left.ptr, sizeof(left.ptr)) );\r
///////////////////////////////////////////////////////////////\r
\r
namespace csbp_krnls\r
-{ \r
- template <int channels> \r
+{\r
+ template <int channels>\r
struct DataCostPerPixel\r
{\r
static __device__ float compute(const uchar* left, const uchar* right)\r
}\r
};\r
\r
- template <> \r
+ template <>\r
struct DataCostPerPixel<1>\r
{\r
static __device__ float compute(const uchar* left, const uchar* right)\r
}\r
};\r
\r
+ template <typename T>\r
+ __global__ void get_first_k_initial_global(T* data_cost_selected_, T *selected_disp_pyr, int h, int w, int nr_plane)\r
+ {\r
+ int x = blockIdx.x * blockDim.x + threadIdx.x;\r
+ int y = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+ if (y < h && x < w)\r
+ {\r
+ T* selected_disparity = selected_disp_pyr + y * cmsg_step1 + x;\r
+ T* data_cost_selected = data_cost_selected_ + y * cmsg_step1 + x;\r
+ T* data_cost = (T*)ctemp + y * cmsg_step1 + x;\r
+\r
+ for(int i = 0; i < nr_plane; i++)\r
+ {\r
+ T fmin_ = data_cost[i * cdisp_step1];\r
+ int id = i;\r
+ for(int j = 0; j < nr_plane; j++)\r
+ {\r
+ T cur = data_cost[j * cdisp_step1];\r
+ if(cur < fmin_)\r
+ {\r
+ fmin_ = cur;\r
+ id = j;\r
+ }\r
+ }\r
+\r
+ data_cost_selected[i * cdisp_step1] = fmin_;\r
+ selected_disparity[i * cdisp_step1] = id;\r
+ data_cost [id * cdisp_step1] = TypeLimits<T>::max();;\r
+ }\r
+ }\r
+ }\r
+\r
+\r
template <typename T>\r
__global__ void get_first_k_initial_local(T* data_cost_selected_, T* selected_disp_pyr, int h, int w, int nr_plane)\r
{\r
int x = blockIdx.x * blockDim.x + threadIdx.x;\r
int y = blockIdx.y * blockDim.y + threadIdx.y;\r
- \r
+\r
if (y < h && x < w)\r
{\r
T* selected_disparity = selected_disp_pyr + y * cmsg_step1 + x;\r
{\r
data_cost_selected[nr_local_minimum * cdisp_step1] = cur;\r
selected_disparity[nr_local_minimum * cdisp_step1] = d;\r
- \r
+\r
data_cost[d * cdisp_step1] = TypeLimits<T>::max();\r
\r
nr_local_minimum++;\r
}\r
\r
template <typename T, int channels>\r
- __global__ void init_data_cost(int h, int w, int level) \r
+ __global__ void init_data_cost(int h, int w, int level)\r
{\r
int x = blockIdx.x * blockDim.x + threadIdx.x;\r
int y = blockIdx.y * blockDim.y + threadIdx.y;\r
- \r
+\r
if (y < h && x < w)\r
{\r
int y0 = y << level;\r
for(int yi = y0; yi < yt; yi++)\r
{\r
for(int xi = x0; xi < xt; xi++)\r
- { \r
+ {\r
int xr = xi - d;\r
- if(d < cth || xr < 0) \r
+ if(d < cth || xr < 0)\r
val += cdata_weight * cmax_data_term;\r
- else \r
- { \r
+ else\r
+ {\r
const uchar* lle = cleft + yi * cimg_step + xi * channels;\r
const uchar* lri = cright + yi * cimg_step + xr * channels;\r
\r
val += DataCostPerPixel<channels>::compute(lle, lri);\r
}\r
- } \r
+ }\r
}\r
data_cost[cdisp_step1 * d] = saturate_cast<T>(val);\r
}\r
}\r
}\r
\r
- template <typename T, int winsz, int channels> \r
+ template <typename T, int winsz, int channels>\r
__global__ void init_data_cost_reduce(int level, int rows, int cols, int h)\r
{\r
- int x_out = blockIdx.x; \r
+ int x_out = blockIdx.x;\r
int y_out = blockIdx.y % h;\r
int d = (blockIdx.y / h) * blockDim.z + threadIdx.z;\r
\r
const uchar* lri = cright + y0 * cimg_step + channels * (x0 + tid - d);\r
\r
for(int y = 0; y < len; ++y)\r
- { \r
+ {\r
val += DataCostPerPixel<channels>::compute(lle, lri);\r
\r
lle += cimg_step;\r
if (winsz >= 32) if (tid < 16) dline[tid] += dline[tid + 16];\r
if (winsz >= 16) if (tid < 8) dline[tid] += dline[tid + 8];\r
if (winsz >= 8) if (tid < 4) dline[tid] += dline[tid + 4];\r
- if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2]; \r
+ if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2];\r
if (winsz >= 2) if (tid < 1) dline[tid] += dline[tid + 1];\r
\r
T* data_cost = (T*)ctemp + y_out * cmsg_step1 + x_out;\r
\r
- if (tid == 0) \r
+ if (tid == 0)\r
data_cost[cdisp_step1 * d] = saturate_cast<T>(dline[0]);\r
}\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
- template <typename T> \r
+ template <typename T>\r
void init_data_cost_caller_(int /*rows*/, int /*cols*/, int h, int w, int level, int /*ndisp*/, int channels, cudaStream_t stream)\r
{\r
dim3 threads(32, 8, 1);\r
dim3 grid(1, 1, 1);\r
\r
grid.x = divUp(w, threads.x);\r
- grid.y = divUp(h, threads.y); \r
- \r
+ grid.y = divUp(h, threads.y);\r
+\r
switch (channels)\r
{\r
case 1: csbp_krnls::init_data_cost<T, 1><<<grid, threads, 0, stream>>>(h, w, level); break;\r
}\r
}\r
\r
- template <typename T, int winsz> \r
+ template <typename T, int winsz>\r
void init_data_cost_reduce_caller_(int rows, int cols, int h, int w, int level, int ndisp, int channels, cudaStream_t stream)\r
{\r
const int threadsNum = 256;\r
const size_t smem_size = threadsNum * sizeof(float);\r
- \r
+\r
dim3 threads(winsz, 1, threadsNum / winsz);\r
- dim3 grid(w, h, 1); \r
+ dim3 grid(w, h, 1);\r
grid.y *= divUp(ndisp, threads.z);\r
- \r
+\r
switch (channels)\r
{\r
case 1: csbp_krnls::init_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(level, rows, cols, h); break;\r
}\r
\r
template<class T>\r
- void init_data_cost_tmpl(int rows, int cols, T* disp_selected_pyr, T* data_cost_selected,\r
- size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, cudaStream_t stream)\r
- { \r
+ void init_data_cost_tmpl(int rows, int cols, T* disp_selected_pyr, T* data_cost_selected, size_t msg_step,\r
+ int h, int w, int level, int nr_plane, int ndisp, int channels, bool use_local_init_data_cost, cudaStream_t stream)\r
+ {\r
\r
typedef void (*InitDataCostCaller)(int cols, int rows, int w, int h, int level, int ndisp, int channels, cudaStream_t stream);\r
\r
- static const InitDataCostCaller init_data_cost_callers[] = \r
+ static const InitDataCostCaller init_data_cost_callers[] =\r
{\r
- init_data_cost_caller_<T>, init_data_cost_caller_<T>, init_data_cost_reduce_caller_<T, 4>, \r
- init_data_cost_reduce_caller_<T, 8>, init_data_cost_reduce_caller_<T, 16>, init_data_cost_reduce_caller_<T, 32>, \r
+ init_data_cost_caller_<T>, init_data_cost_caller_<T>, init_data_cost_reduce_caller_<T, 4>,\r
+ init_data_cost_reduce_caller_<T, 8>, init_data_cost_reduce_caller_<T, 16>, init_data_cost_reduce_caller_<T, 32>,\r
init_data_cost_reduce_caller_<T, 64>, init_data_cost_reduce_caller_<T, 128>, init_data_cost_reduce_caller_<T, 256>\r
};\r
- \r
+\r
size_t disp_step = msg_step * h;\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step1, &disp_step, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step1, &msg_step, sizeof(size_t)) );\r
grid.x = divUp(w, threads.x);\r
grid.y = divUp(h, threads.y);\r
\r
- csbp_krnls::get_first_k_initial_local<<<grid, threads, 0, stream>>>(data_cost_selected, disp_selected_pyr, h, w, nr_plane); \r
+ if (use_local_init_data_cost == true)\r
+ csbp_krnls::get_first_k_initial_local<<<grid, threads, 0, stream>>> (data_cost_selected, disp_selected_pyr, h, w, nr_plane);\r
+ else\r
+ csbp_krnls::get_first_k_initial_global<<<grid, threads, 0, stream>>>(data_cost_selected, disp_selected_pyr, h, w, nr_plane);\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
void init_data_cost(int rows, int cols, short* disp_selected_pyr, short* data_cost_selected,\r
- size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, cudaStream_t stream)\r
+ 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)\r
{\r
- init_data_cost_tmpl(rows, cols, disp_selected_pyr, data_cost_selected, msg_step, h, w, level, nr_plane, ndisp, channels, stream);\r
+ init_data_cost_tmpl(rows, cols, disp_selected_pyr, data_cost_selected, msg_step, h, w, level, nr_plane, ndisp, channels, use_local_init_data_cost, stream);\r
}\r
\r
void init_data_cost(int rows, int cols, float* disp_selected_pyr, float* data_cost_selected,\r
- size_t msg_step, int h, int w, int level, int nr_plane, int ndisp, int channels, cudaStream_t stream)\r
+ 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)\r
{\r
- init_data_cost_tmpl(rows, cols, disp_selected_pyr, data_cost_selected, msg_step, h, w, level, nr_plane, ndisp, channels, stream);\r
+ init_data_cost_tmpl(rows, cols, disp_selected_pyr, data_cost_selected, msg_step, h, w, level, nr_plane, ndisp, channels, use_local_init_data_cost, stream);\r
}\r
\r
}}}\r
__global__ void compute_data_cost(const T* selected_disp_pyr, T* data_cost_, int h, int w, int level, int nr_plane)\r
{\r
int x = blockIdx.x * blockDim.x + threadIdx.x;\r
- int y = blockIdx.y * blockDim.y + threadIdx.y; \r
+ int y = blockIdx.y * blockDim.y + threadIdx.y;\r
\r
if (y < h && x < w)\r
{\r
int y0 = y << level;\r
int yt = (y + 1) << level;\r
- \r
+\r
int x0 = x << level;\r
int xt = (x + 1) << level;\r
\r
int sel_disp = selected_disparity[d * cdisp_step2];\r
int xr = xi - sel_disp;\r
\r
- if (xr < 0 || sel_disp < cth) \r
+ if (xr < 0 || sel_disp < cth)\r
val += cdata_weight * cmax_data_term;\r
- else \r
+ else\r
{\r
const uchar* left_x = cleft + yi * cimg_step + xi * channels;\r
const uchar* right_x = cright + yi * cimg_step + xr * channels;\r
}\r
}\r
\r
- template <typename T, int winsz, int channels> \r
+ template <typename T, int winsz, int channels>\r
__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)\r
{\r
- int x_out = blockIdx.x; \r
+ int x_out = blockIdx.x;\r
int y_out = blockIdx.y % h;\r
int d = (blockIdx.y / h) * blockDim.z + threadIdx.z;\r
\r
int tid = threadIdx.x;\r
\r
const T* selected_disparity = selected_disp_pyr + y_out/2 * cmsg_step2 + x_out/2;\r
- T* data_cost = data_cost_ + y_out * cmsg_step1 + x_out; \r
+ T* data_cost = data_cost_ + y_out * cmsg_step1 + x_out;\r
\r
if (d < nr_plane)\r
{\r
const uchar* lri = cright + y0 * cimg_step + channels * (x0 + tid - sel_disp);\r
\r
for(int y = 0; y < len; ++y)\r
- { \r
+ {\r
val += DataCostPerPixel<channels>::compute(lle, lri);\r
\r
lle += cimg_step;\r
if (winsz >= 32) if (tid < 16) dline[tid] += dline[tid + 16];\r
if (winsz >= 16) if (tid < 8) dline[tid] += dline[tid + 8];\r
if (winsz >= 8) if (tid < 4) dline[tid] += dline[tid + 4];\r
- if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2]; \r
+ if (winsz >= 4) if (tid < 2) dline[tid] += dline[tid + 2];\r
if (winsz >= 2) if (tid < 1) dline[tid] += dline[tid + 1];\r
\r
- if (tid == 0) \r
+ if (tid == 0)\r
data_cost[cdisp_step1 * d] = saturate_cast<T>(dline[0]);\r
}\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
- template <typename T> \r
+ template <typename T>\r
void compute_data_cost_caller_(const T* disp_selected_pyr, T* data_cost, int /*rows*/, int /*cols*/,\r
int h, int w, int level, int nr_plane, int channels, cudaStream_t stream)\r
{\r
case 1: csbp_krnls::compute_data_cost<T, 1><<<grid, threads, 0, stream>>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break;\r
case 3: csbp_krnls::compute_data_cost<T, 3><<<grid, threads, 0, stream>>>(disp_selected_pyr, data_cost, h, w, level, nr_plane); break;\r
default: cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- } \r
+ }\r
}\r
\r
- template <typename T, int winsz> \r
+ template <typename T, int winsz>\r
void compute_data_cost_reduce_caller_(const T* disp_selected_pyr, T* data_cost, int rows, int cols,\r
int h, int w, int level, int nr_plane, int channels, cudaStream_t stream)\r
{\r
const int threadsNum = 256;\r
const size_t smem_size = threadsNum * sizeof(float);\r
- \r
+\r
dim3 threads(winsz, 1, threadsNum / winsz);\r
- dim3 grid(w, h, 1); \r
+ dim3 grid(w, h, 1);\r
grid.y *= divUp(nr_plane, threads.z);\r
- \r
+\r
switch (channels)\r
{\r
case 1: csbp_krnls::compute_data_cost_reduce<T, winsz, 1><<<grid, threads, smem_size, stream>>>(disp_selected_pyr, data_cost, level, rows, cols, h, nr_plane); break;\r
default: cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
}\r
}\r
- \r
- \r
+\r
+\r
template<class T>\r
void compute_data_cost_tmpl(const T* disp_selected_pyr, T* data_cost, size_t msg_step1, size_t msg_step2,\r
int rows, int cols, int h, int w, int h2, int level, int nr_plane, int channels, cudaStream_t stream)\r
{\r
- typedef void (*ComputeDataCostCaller)(const T* disp_selected_pyr, T* data_cost, int rows, int cols, \r
+ typedef void (*ComputeDataCostCaller)(const T* disp_selected_pyr, T* data_cost, int rows, int cols,\r
int h, int w, int level, int nr_plane, int channels, cudaStream_t stream);\r
\r
- static const ComputeDataCostCaller callers[] = \r
+ static const ComputeDataCostCaller callers[] =\r
{\r
- compute_data_cost_caller_<T>, compute_data_cost_caller_<T>, compute_data_cost_reduce_caller_<T, 4>, \r
- compute_data_cost_reduce_caller_<T, 8>, compute_data_cost_reduce_caller_<T, 16>, compute_data_cost_reduce_caller_<T, 32>, \r
+ compute_data_cost_caller_<T>, compute_data_cost_caller_<T>, compute_data_cost_reduce_caller_<T, 4>,\r
+ compute_data_cost_reduce_caller_<T, 8>, compute_data_cost_reduce_caller_<T, 16>, compute_data_cost_reduce_caller_<T, 32>,\r
compute_data_cost_reduce_caller_<T, 64>, compute_data_cost_reduce_caller_<T, 128>, compute_data_cost_reduce_caller_<T, 256>\r
};\r
\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step1, &disp_step1, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step2, &disp_step2, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step1, &msg_step1, sizeof(size_t)) );\r
- cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step2, &msg_step2, sizeof(size_t)) ); \r
- \r
+ cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step2, &msg_step2, sizeof(size_t)) );\r
+\r
callers[level](disp_selected_pyr, data_cost, rows, cols, h, w, level, nr_plane, channels, stream);\r
- \r
+\r
if (stream == 0)\r
- cudaSafeCall( cudaThreadSynchronize() ); \r
+ cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
void compute_data_cost(const short* disp_selected_pyr, short* data_cost, size_t msg_step1, size_t msg_step2,\r
namespace csbp_krnls\r
{\r
template <typename T>\r
- __device__ void get_first_k_element_increase(T* u_new, T* d_new, T* l_new, T* r_new, \r
+ __device__ void get_first_k_element_increase(T* u_new, T* d_new, T* l_new, T* r_new,\r
const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur,\r
- T* data_cost_selected, T* disparity_selected_new, T* data_cost_new, \r
- const T* data_cost_cur, const T* disparity_selected_cur, \r
+ T* data_cost_selected, T* disparity_selected_new, T* data_cost_new,\r
+ const T* data_cost_cur, const T* disparity_selected_cur,\r
int nr_plane, int nr_plane2)\r
{\r
for(int i = 0; i < nr_plane; i++)\r
}\r
\r
template <typename T>\r
- __global__ void init_message(T* u_new_, T* d_new_, T* l_new_, T* r_new_, \r
- const T* u_cur_, const T* d_cur_, const T* l_cur_, const T* r_cur_, \r
- T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, \r
- T* data_cost_selected_, const T* data_cost_, \r
+ __global__ void init_message(T* u_new_, T* d_new_, T* l_new_, T* r_new_,\r
+ const T* u_cur_, const T* d_cur_, const T* l_cur_, const T* r_cur_,\r
+ T* selected_disp_pyr_new, const T* selected_disp_pyr_cur,\r
+ T* data_cost_selected_, const T* data_cost_,\r
int h, int w, int nr_plane, int h2, int w2, int nr_plane2)\r
{\r
int x = blockIdx.x * blockDim.x + threadIdx.x;\r
int y = blockIdx.y * blockDim.y + threadIdx.y;\r
\r
if (y < h && x < w)\r
- { \r
+ {\r
const T* u_cur = u_cur_ + min(h2-1, y/2 + 1) * cmsg_step2 + x/2;\r
const T* d_cur = d_cur_ + max(0, y/2 - 1) * cmsg_step2 + x/2;\r
const T* l_cur = l_cur_ + y/2 * cmsg_step2 + min(w2-1, x/2 + 1);\r
for(int d = 0; d < nr_plane2; d++)\r
{\r
int idx2 = d * cdisp_step2;\r
- \r
+\r
T val = data_cost[d * cdisp_step1] + u_cur[idx2] + d_cur[idx2] + l_cur[idx2] + r_cur[idx2];\r
data_cost_new[d * cdisp_step1] = val;\r
}\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
template<class T>\r
- void init_message_tmpl(T* u_new, T* d_new, T* l_new, T* r_new, \r
- const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur, \r
- T* selected_disp_pyr_new, const T* selected_disp_pyr_cur, \r
- T* data_cost_selected, const T* data_cost, size_t msg_step1, size_t msg_step2, \r
+ void init_message_tmpl(T* u_new, T* d_new, T* l_new, T* r_new,\r
+ const T* u_cur, const T* d_cur, const T* l_cur, const T* r_cur,\r
+ T* selected_disp_pyr_new, const T* selected_disp_pyr_cur,\r
+ T* data_cost_selected, const T* data_cost, size_t msg_step1, size_t msg_step2,\r
int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream)\r
- { \r
- \r
+ {\r
+\r
size_t disp_step1 = msg_step1 * h;\r
size_t disp_step2 = msg_step2 * h2;\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step1, &disp_step1, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step2, &disp_step2, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step1, &msg_step1, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step2, &msg_step2, sizeof(size_t)) );\r
- \r
+\r
dim3 threads(32, 8, 1);\r
dim3 grid(1, 1, 1);\r
\r
grid.x = divUp(w, threads.x);\r
- grid.y = divUp(h, threads.y); \r
+ grid.y = divUp(h, threads.y);\r
\r
- csbp_krnls::init_message<<<grid, threads, 0, stream>>>(u_new, d_new, l_new, r_new, \r
+ csbp_krnls::init_message<<<grid, threads, 0, stream>>>(u_new, d_new, l_new, r_new,\r
u_cur, d_cur, l_cur, r_cur,\r
- selected_disp_pyr_new, selected_disp_pyr_cur, \r
- data_cost_selected, data_cost, \r
+ selected_disp_pyr_new, selected_disp_pyr_cur,\r
+ data_cost_selected, data_cost,\r
h, w, nr_plane, h2, w2, nr_plane2);\r
- \r
+\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- void init_message(short* u_new, short* d_new, short* l_new, short* r_new, \r
- const short* u_cur, const short* d_cur, const short* l_cur, const short* r_cur, \r
- short* selected_disp_pyr_new, const short* selected_disp_pyr_cur, \r
- short* data_cost_selected, const short* data_cost, size_t msg_step1, size_t msg_step2, \r
+ void init_message(short* u_new, short* d_new, short* l_new, short* r_new,\r
+ const short* u_cur, const short* d_cur, const short* l_cur, const short* r_cur,\r
+ short* selected_disp_pyr_new, const short* selected_disp_pyr_cur,\r
+ short* data_cost_selected, const short* data_cost, size_t msg_step1, size_t msg_step2,\r
int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream)\r
{\r
- init_message_tmpl(u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur, \r
- selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, msg_step1, msg_step2, \r
+ init_message_tmpl(u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur,\r
+ selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, msg_step1, msg_step2,\r
h, w, nr_plane, h2, w2, nr_plane2, stream);\r
}\r
\r
- void init_message(float* u_new, float* d_new, float* l_new, float* r_new, \r
- const float* u_cur, const float* d_cur, const float* l_cur, const float* r_cur, \r
- float* selected_disp_pyr_new, const float* selected_disp_pyr_cur, \r
- float* data_cost_selected, const float* data_cost, size_t msg_step1, size_t msg_step2, \r
+ void init_message(float* u_new, float* d_new, float* l_new, float* r_new,\r
+ const float* u_cur, const float* d_cur, const float* l_cur, const float* r_cur,\r
+ float* selected_disp_pyr_new, const float* selected_disp_pyr_cur,\r
+ float* data_cost_selected, const float* data_cost, size_t msg_step1, size_t msg_step2,\r
int h, int w, int nr_plane, int h2, int w2, int nr_plane2, cudaStream_t stream)\r
{\r
- init_message_tmpl(u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur, \r
- selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, msg_step1, msg_step2, \r
+ init_message_tmpl(u_new, d_new, l_new, r_new, u_cur, d_cur, l_cur, r_cur,\r
+ selected_disp_pyr_new, selected_disp_pyr_cur, data_cost_selected, data_cost, msg_step1, msg_step2,\r
h, w, nr_plane, h2, w2, nr_plane2, stream);\r
}\r
}}}\r
namespace csbp_krnls\r
{\r
template <typename T>\r
- __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3, \r
+ __device__ void message_per_pixel(const T* data, T* msg_dst, const T* msg1, const T* msg2, const T* msg3,\r
const T* dst_disp, const T* src_disp, int nr_plane, T* temp)\r
{\r
T minimum = TypeLimits<T>::max();\r
int idx = d * cdisp_step1;\r
T val = data[idx] + msg1[idx] + msg2[idx] + msg3[idx];\r
\r
- if(val < minimum) \r
+ if(val < minimum)\r
minimum = val;\r
\r
msg_dst[idx] = val;\r
\r
for(int d2 = 0; d2 < nr_plane; d2++)\r
cost_min = fmin(cost_min, msg_dst[d2 * cdisp_step1] + cdisc_single_jump * abs(dst_disp[d2 * cdisp_step1] - src_disp_reg));\r
- \r
+\r
temp[d * cdisp_step1] = saturate_cast<T>(cost_min);\r
sum += cost_min;\r
}\r
T* d = d_ + y * cmsg_step1 + x;\r
T* l = l_ + y * cmsg_step1 + x;\r
T* r = r_ + y * cmsg_step1 + x;\r
- \r
+\r
const T* disp = selected_disp_pyr_cur + y * cmsg_step1 + x;\r
- \r
+\r
T* temp = (T*)ctemp + y * cmsg_step1 + x;\r
\r
message_per_pixel(data, u, r - 1, u + cmsg_step1, l + 1, disp, disp - cmsg_step1, nr_plane, temp);\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
template<class T>\r
- void calc_all_iterations_tmpl(T* u, T* d, T* l, T* r, const T* data_cost_selected, \r
+ void calc_all_iterations_tmpl(T* u, T* d, T* l, T* r, const T* data_cost_selected,\r
const T* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream)\r
- { \r
+ {\r
size_t disp_step = msg_step * h;\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step1, &disp_step, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step1, &msg_step, sizeof(size_t)) );\r
\r
for(int t = 0; t < iters; ++t)\r
{\r
- csbp_krnls::compute_message<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1); \r
- \r
+ csbp_krnls::compute_message<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, h, w, nr_plane, t & 1);\r
+\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
};\r
\r
- void calc_all_iterations(short* u, short* d, short* l, short* r, short* data_cost_selected, \r
- const short* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream) \r
+ void calc_all_iterations(short* u, short* d, short* l, short* r, short* data_cost_selected,\r
+ const short* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream)\r
{\r
calc_all_iterations_tmpl(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, msg_step, h, w, nr_plane, iters, stream);\r
}\r
\r
- void calc_all_iterations(float*u, float* d, float* l, float* r, float* data_cost_selected, \r
+ void calc_all_iterations(float*u, float* d, float* l, float* r, float* data_cost_selected,\r
const float* selected_disp_pyr_cur, size_t msg_step, int h, int w, int nr_plane, int iters, cudaStream_t stream)\r
{\r
calc_all_iterations_tmpl(u, d, l, r, data_cost_selected, selected_disp_pyr_cur, msg_step, h, w, nr_plane, iters, stream);\r
namespace csbp_krnls\r
{\r
template <typename T>\r
- __global__ void compute_disp(const T* u_, const T* d_, const T* l_, const T* r_, \r
- const T* data_cost_selected, const T* disp_selected_pyr, \r
- short* disp, size_t res_step, int cols, int rows, int nr_plane) \r
- { \r
+ __global__ void compute_disp(const T* u_, const T* d_, const T* l_, const T* r_,\r
+ const T* data_cost_selected, const T* disp_selected_pyr,\r
+ short* disp, size_t res_step, int cols, int rows, int nr_plane)\r
+ {\r
int x = blockIdx.x * blockDim.x + threadIdx.x;\r
int y = blockIdx.y * blockDim.y + threadIdx.y;\r
\r
const T* d = d_ + (y-1) * cmsg_step1 + (x+0);\r
const T* l = l_ + (y+0) * cmsg_step1 + (x+1);\r
const T* r = r_ + (y+0) * cmsg_step1 + (x-1);\r
- \r
+\r
int best = 0;\r
T best_val = TypeLimits<T>::max();\r
- for (int i = 0; i < nr_plane; ++i) \r
+ for (int i = 0; i < nr_plane; ++i)\r
{\r
int idx = i * cdisp_step1;\r
T val = data[idx]+ u[idx] + d[idx] + l[idx] + r[idx];\r
\r
- if (val < best_val) \r
+ if (val < best_val)\r
{\r
best_val = val;\r
best = saturate_cast<short>(disp_selected[idx]);\r
}\r
}\r
\r
-namespace cv { namespace gpu { namespace csbp \r
+namespace cv { namespace gpu { namespace csbp\r
{\r
- template<class T> \r
- void compute_disp_tmpl(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, \r
+ template<class T>\r
+ void compute_disp_tmpl(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,\r
const DevMem2D_<short>& disp, int nr_plane, cudaStream_t stream)\r
- { \r
+ {\r
size_t disp_step = disp.rows * msg_step;\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cdisp_step1, &disp_step, sizeof(size_t)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_krnls::cmsg_step1, &msg_step, sizeof(size_t)) );\r
\r
grid.x = divUp(disp.cols, threads.x);\r
grid.y = divUp(disp.rows, threads.y);\r
- \r
- csbp_krnls::compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected, \r
+\r
+ csbp_krnls::compute_disp<<<grid, threads, 0, stream>>>(u, d, l, r, data_cost_selected, disp_selected,\r
disp.ptr, disp.step / disp.elemSize(), disp.cols, disp.rows, nr_plane);\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- void compute_disp(const short* u, const short* d, const short* l, const short* r, const short* data_cost_selected, const short* disp_selected, size_t msg_step, \r
+ void compute_disp(const short* u, const short* d, const short* l, const short* r, const short* data_cost_selected, const short* disp_selected, size_t msg_step,\r
DevMem2D_<short> disp, int nr_plane, cudaStream_t stream)\r
{\r
compute_disp_tmpl(u, d, l, r, data_cost_selected, disp_selected, msg_step, disp, nr_plane, stream);\r
}\r
\r
- void compute_disp(const float* u, const float* d, const float* l, const float* r, const float* data_cost_selected, const float* disp_selected, size_t msg_step, \r
+ void compute_disp(const float* u, const float* d, const float* l, const float* r, const float* data_cost_selected, const float* disp_selected, size_t msg_step,\r
DevMem2D_<short> disp, int nr_plane, cudaStream_t stream)\r
{\r
compute_disp_tmpl(u, d, l, r, data_cost_selected, disp_selected, msg_step, disp, nr_plane, stream);\r
}\r
\r
-}}}
\ No newline at end of file
+}}}\r
projects / platform / upstream / opencv.git / commitdiff