GpuMat& buf, Stream& stream = Stream::Null());\r
\r
//! performs connected componnents labeling.\r
-CV_EXPORTS void labelComponents(const GpuMat& image, GpuMat& mask, GpuMat& components, const cv::Scalar& lo, const cv::Scalar& hi);\r
+CV_EXPORTS void labelComponents(const GpuMat& image, GpuMat& mask, GpuMat& components, const cv::Scalar& lo, const cv::Scalar& hi, Stream& stream = Stream::Null());\r
\r
////////////////////////////////// Histograms //////////////////////////////////\r
\r
};
template<typename F>
- __global__ void computeComponents(const DevMem2D image, DevMem2D components, F connected)
+ __global__ void computeConnectivity(const DevMem2D image, DevMem2D components, F connected)
{
int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;
components(y, x) = c;
}
- void computeEdges(const DevMem2D& image, DevMem2D components, const int lo, const int hi)
+ void computeEdges(const DevMem2D& image, DevMem2D edges, const int lo, const int hi, cudaStream_t stream)
{
dim3 block(CTA_SIZE_X, CTA_SIZE_Y);
dim3 grid(divUp(image.cols, block.x), divUp(image.rows, block.y));
InInterval<int> inInt(lo, hi);
- computeComponents<InInterval<int> ><<<grid, block>>>(image, components, inInt);
+ computeConnectivity<InInterval<int> ><<<grid, block, 0, stream>>>(image, edges, inInt);
cudaSafeCall( cudaGetLastError() );
- cudaSafeCall( cudaDeviceSynchronize() );
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
}
__global__ void lableTiles(const DevMem2D edges, DevMem2Di comps)
comps(y, x) = root(comps, comps(y, x));
}
- void labelComponents(const DevMem2D& edges, DevMem2Di comps)
+ void labelComponents(const DevMem2D& edges, DevMem2Di comps, cudaStream_t stream)
{
dim3 block(CTA_SIZE_X, CTA_SIZE_Y);
dim3 grid(divUp(edges.cols, TILE_COLS), divUp(edges.rows, TILE_ROWS));
- lableTiles<<<grid, block>>>(edges, comps);
+ lableTiles<<<grid, block, 0, stream>>>(edges, comps);
cudaSafeCall( cudaGetLastError() );
int tileSizeX = TILE_COLS, tileSizeY = TILE_ROWS;
cudaSafeCall( cudaGetLastError() );
- cudaSafeCall( cudaDeviceSynchronize() );
+ // cudaSafeCall( cudaDeviceSynchronize() );
while (grid.x > 1 || grid.y > 1)
{
dim3 mergeGrid(ceilf(grid.x / 2.0), ceilf(grid.y / 2.0));
dim3 mergeBlock(STA_SIZE_MARGE_X, STA_SIZE_MARGE_Y);
std::cout << "merging: " << grid.y << " x " << grid.x << " ---> " << mergeGrid.y << " x " << mergeGrid.x << " for tiles: " << tileSizeY << " x " << tileSizeX << std::endl;
- crossMerge<<<mergeGrid, mergeBlock>>>(2, 2, tileSizeY, tileSizeX, edges, comps, ceilf(grid.y / 2.0) - grid.y / 2, ceilf(grid.x / 2.0) - grid.x / 2);
+ crossMerge<<<mergeGrid, mergeBlock, 0, stream>>>(2, 2, tileSizeY, tileSizeX, edges, comps, ceilf(grid.y / 2.0) - grid.y / 2, ceilf(grid.x / 2.0) - grid.x / 2);
tileSizeX <<= 1;
tileSizeY <<= 1;
grid = mergeGrid;
grid.x = divUp(edges.cols, block.x);
grid.y = divUp(edges.rows, block.y);
- flatten<<<grid, block>>>(edges, comps);
+ flatten<<<grid, block, 0, stream>>>(edges, comps);
cudaSafeCall( cudaGetLastError() );
- cudaSafeCall( cudaDeviceSynchronize() );
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
}
}
} } }
\ No newline at end of file
void cv::gpu::graphcut(GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
void cv::gpu::graphcut(GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
\r
-void cv::gpu::labelComponents(const GpuMat& image, GpuMat& mask, GpuMat& components, const cv::Scalar& lo, const cv::Scalar& hi) { throw_nogpu(); }\r
+void cv::gpu::labelComponents(const GpuMat&, GpuMat&, GpuMat&, const cv::Scalar&, const cv::Scalar&, Stream&) { throw_nogpu(); }\r
\r
#else /* !defined (HAVE_CUDA) */\r
\r
{\r
namespace ccl\r
{\r
- void labelComponents(const DevMem2D& edges, DevMem2Di comps);\r
- void computeEdges(const DevMem2D& image, DevMem2D edges, const int lo, const int hi);\r
+ void labelComponents(const DevMem2D& edges, DevMem2Di comps, cudaStream_t stream);\r
+ void computeEdges(const DevMem2D& image, DevMem2D edges, const int lo, const int hi, cudaStream_t stream);\r
}\r
}}}\r
\r
-void cv::gpu::labelComponents(const GpuMat& image, GpuMat& mask, GpuMat& components, const cv::Scalar& lo, const cv::Scalar& hi)\r
+void cv::gpu::labelComponents(const GpuMat& image, GpuMat& mask, GpuMat& components, const cv::Scalar& lo, const cv::Scalar& hi, Stream& s)\r
{\r
- device::ccl::computeEdges(image, mask, lo[0], hi[0]);\r
- device::ccl::labelComponents(mask, components);\r
+ CV_Assert(!image.empty());\r
+\r
+ int type = image.type();\r
+ CV_Assert(type == CV_8UC1);\r
+\r
+ if (image.size() != mask.size() || mask.type() != CV_8UC1)\r
+ mask.create(image.size(), CV_8UC1);\r
+\r
+ if (image.size() != components.size() || components.type() != CV_32SC1)\r
+ components.create(image.size(), CV_32SC1);\r
+\r
+ cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
+ device::ccl::computeEdges(image, mask, lo[0], hi[0], stream);\r
+ device::ccl::labelComponents(mask, components, stream);\r
}\r
\r
\r
{
cv::Mat image;
cvtColor(loat_image(), image, CV_BGR2GRAY);
- cv::Mat image_cpu = image.clone();
-
- // cv::floodFill(image, cv::Point(1,1),cv::Scalar::all(64), 0, cv::Scalar::all(0), cv::Scalar::all(256));
cv::gpu::GpuMat mask;
mask.create(image.rows, image.cols, CV_8UC1);
cv::gpu::GpuMat components;
components.create(image.rows, image.cols, CV_32SC1);
- std::cout << "summary: " << image.cols << " " << image.rows << " "
- << cv::gpu::GpuMat(image).cols << " " << cv::gpu::GpuMat(image).rows<< " "
- << mask.cols << " " << mask.rows<< " "
- << components.cols << " " << components.rows<< std::endl;
-
-
cv::gpu::labelComponents(cv::gpu::GpuMat(image), mask, components, cv::Scalar::all(0), cv::Scalar::all(2));
- // // for(int i = 0; i + 32 < image.rows; i += 32)
- // // for(int j = 0; j + 32 < image.cols; j += 32)
- // // {
- // // std::cout << cv::Mat(cv::Mat(mask), cv::Rect(j, i, 32, 32 ))<< std::endl;
- // // std::cout << cv::Mat(cv::Mat(components), cv::Rect(j, i, 32, 32 )) << std::endl;
- // // }
-
// std::cout << cv::Mat(components) << std::endl;
// cv::imshow("test", image);
// cv::waitKey(0);
cv::waitKey(0);
cv::imshow("test", cv::Mat(components) * 2);
cv::waitKey(0);
- std::cout << "test! " << image.cols << std::endl;
}
INSTANTIATE_TEST_CASE_P(ConnectedComponents, Labeling, ALL_DEVICES);
\ No newline at end of file