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

cuda::DisparityBilateralFilter no longer uses constant memory for parameters

Browse Source 
Now multiple filters can be used in the same context without stepping on each other.
This commit is contained in:
Aaron Denney 2014-06-23 13:55:09 -07:00
parent 8f8450793a
commit d848704b35
2 changed files with 15 additions and 37 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -49,30 +49,6 @@ namespace cv { namespace cuda { namespace device
{ {
namespace disp_bilateral_filter 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> template <int channels>
struct DistRgbMax struct DistRgbMax
{ {
@ -95,7 +71,11 @@ namespace cv { namespace cuda { namespace device
}; };
template <int channels, typename T> 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 y = blockIdx.y * blockDim.y + threadIdx.y;
const int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + t) & 1); const int x = ((blockIdx.x * blockDim.x + threadIdx.x) << 1) + ((y + t) & 1);
@ -178,7 +158,7 @@ namespace cv { namespace cuda { namespace device
} }
template <typename T> 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 threads(32, 8, 1);
dim3 grid(1, 1, 1); dim3 grid(1, 1, 1);
@ -190,20 +170,20 @@ namespace cv { namespace cuda { namespace device
case 1: case 1:
for (int i = 0; i < iters; ++i) 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() ); 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() ); cudaSafeCall( cudaGetLastError() );
} }
break; break;
case 3: case 3:
for (int i = 0; i < iters; ++i) 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() ); 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() ); cudaSafeCall( cudaGetLastError() );
} }
break; break;
@ -215,8 +195,8 @@ namespace cv { namespace cuda { namespace device
cudaSafeCall( cudaDeviceSynchronize() ); 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<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, 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 bilateral_filter
}}} // namespace cv { namespace cuda { namespace cudev }}} // namespace cv { namespace cuda { namespace cudev

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

@ -55,10 +55,8 @@ namespace cv { namespace cuda { namespace device
{ {
namespace disp_bilateral_filter 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> 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 *, const float *, size_t, int radius, short edge_disc, short max_disc, cudaStream_t stream);
} }
}}} }}}
@ -165,7 +163,7 @@ namespace
const short edge_disc = std::max<short>(short(1), short(ndisp * edge_threshold + 0.5)); 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); 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()); _dst.create(disp.size(), disp.type());
GpuMat dst = _dst.getGpuMat(); GpuMat dst = _dst.getGpuMat();
@ -173,7 +171,7 @@ namespace
if (dst.data != disp.data) if (dst.data != disp.data)
disp.copyTo(dst, stream); 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) void DispBilateralFilterImpl::apply(InputArray _disp, InputArray _image, OutputArray dst, Stream& stream)